ComputePrimitives.h

Go to the documentation of this file.

// Copyright Contributors to the OpenVDB Project
// SPDX-License-Identifier: Apache-2.0
 
/// @file ComputePrimitives.h
/// @brief A collection of parallel compute primitives
 
#pragma once
 
#if defined(NANOVDB_USE_CUDA)
#include <cuda_runtime_api.h>
#endif
 
#if defined(NANOVDB_USE_TBB)
#include <tbb/parallel_for.h>
#include <tbb/blocked_range.h>
#endif
 
#include <utility>
#include <tuple>
 
 
#if defined(__CUDACC__)
 
static inline bool checkCUDA(cudaError_t result, const char* file, const int line)
{
    if (result != cudaSuccess) {
        std::cerr << "CUDA Runtime API error " << result << " in file " << file << ", line " << line << " : " << cudaGetErrorString(result) << ".\n";
        return false;
    }
    return true;
}
 
#define NANOVDB_CUDA_SAFE_CALL(x) checkCUDA(x, __FILE__, __LINE__)
 
static inline void checkErrorCUDA(cudaError_t result, const char* file, const int line)
{
    if (result != cudaSuccess) {
        std::cerr << "CUDA Runtime API error " << result << " in file " << file << ", line " << line << " : " << cudaGetErrorString(result) << ".\n";
        exit(1);
    }
}
 
#define NANOVDB_CUDA_CHECK_ERROR(result, file, line) checkErrorCUDA(result, file, line)
 
#endif
 
template<typename Fn, typename... Args>
class ApplyFunc
{
public:
    ApplyFunc(int count, int blockSize, const Fn& fn, Args... args)
        : mCount(count)
        , mBlockSize(blockSize)
        , mArgs(args...)
        , mFunc(fn)
    {
    }
 
    template<std::size_t... Is>
    void call(int start, int end, std::index_sequence<Is...>) const
    {
        mFunc(start, end, std::get<Is>(mArgs)...);
    }
 
    void operator()(int i) const
    {
        int start = i * mBlockSize;
        int end = i * mBlockSize + mBlockSize;
        if (end > mCount)
            end = mCount;
        call(start, end, std::make_index_sequence<sizeof...(Args)>());
    }
 
#if defined(NANOVDB_USE_TBB)
    void operator()(const tbb::blocked_range<int>& r) const
    {
        int start = r.begin();
        int end = r.end();
        if (end > mCount)
            end = mCount;
        call(start, end, std::make_index_sequence<sizeof...(Args)>());
    }
#endif
 
private:
    int                 mCount;
    int                 mBlockSize;
    Fn                  mFunc;
    std::tuple<Args...> mArgs;
};
 
#if defined(__CUDACC__)
 
template<int WorkPerThread, typename FnT, typename... Args>
__global__ void parallelForKernel(int numItems, FnT f, Args... args)
{
    for (int j=0;j<WorkPerThread;++j)
    {
        int i = threadIdx.x + blockIdx.x * blockDim.x + j * blockDim.x * gridDim.x;
        if (i < numItems)
            f(i, i + 1, args...);
    }
}
 
#endif
 
inline void computeSync(bool useCuda, const char* file, int line)
{
#if defined(__CUDACC__)
    if (useCuda) {
        NANOVDB_CUDA_CHECK_ERROR(cudaDeviceSynchronize(), file, line);
    }
#endif
}
 
inline void computeFill(bool useCuda, void* data, uint8_t value, size_t size)
{
    if (useCuda) {
#if defined(__CUDACC__)
        cudaMemset(data, value, size);
#endif
    } else {
        std::memset(data, value, size);
    }
}
 
template<typename FunctorT, typename... Args>
inline void computeForEach(bool useCuda, int numItems, int blockSize, const char* file, int line, const FunctorT& op, Args... args)
{
    if (numItems == 0)
        return;
 
    if (useCuda) {
#if defined(__CUDACC__)
        static const int WorkPerThread = 1;
        int blockCount = ((numItems/WorkPerThread) + (blockSize - 1)) / blockSize;
        parallelForKernel<WorkPerThread, FunctorT, Args...><<<blockCount, blockSize, 0, 0>>>(numItems, op, args...);
        NANOVDB_CUDA_CHECK_ERROR(cudaGetLastError(), file, line);
#endif
    } else {
#if defined(NANOVDB_USE_TBB)
        tbb::blocked_range<int> range(0, numItems, blockSize);
        tbb::parallel_for(range, ApplyFunc<FunctorT, Args...>(numItems, blockSize, op, args...));
#else
        for (int i = 0; i < numItems; ++i)
            op(i, i + 1, args...);
#endif
    }
}
 
inline void computeDownload(bool useCuda, void* dst, const void* src, size_t size)
{
    if (useCuda) {
#if defined(__CUDACC__)
        cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost);
#endif
    } else {
        std::memcpy(dst, src, size);
    }
}
 
inline void computeCopy(bool useCuda, void* dst, const void* src, size_t size)
{
    if (useCuda) {
#if defined(__CUDACC__)
        cudaMemcpy(dst, src, size, cudaMemcpyDeviceToDevice);
#endif
    } else {
        std::memcpy(dst, src, size);
    }
}