cuda_utils.cuh

Go to the documentation of this file.

 
#pragma once
#include <bit> // bit_width
#include <cstddef> // size_t
#include <cstdint> // uintmax_t
#include <cstdio> // fprintf
#include <cuda_runtime.h> // CUDA runtime
#include <limits> // std::numeric_limits
 
#define BLOCKS 64
#define THREADS 1024
 
#define gpuErrchk(ans) { gpuAssert(ans, __FILE__, __LINE__); }
 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
 
        if (abort)
        {
            exit(code);
        }
    }
}
 
static constexpr bool CHECK_CUDA_ERRORS = false;
 
__host__ inline size_t host_fast_log2(size_t n)
{
    // std::bit_width(n) returns number of bits needed to represent n.
    // If n is power of 2, say 8 (1000), bit_width is 4. log2(8) = 3. So bit_width - 1.
    // If n is not power of 2, say 7 (0111), bit_width is 3. log2(7) is approx 2.8. floor(log2(7)) = 2. So bit_width - 1.
    return std::bit_width(n) - 1;
}
 
__host__ inline size_t host_fast_log2_ceil(size_t n)
{
    // std::bit_width(n-1) returns number of bits needed for n-1.
    // If n is power of 2, say 8 (1000), n-1 is 7 (0111). bit_width(7) = 3. ceil(log2(8)) = 3.
    // If n is not power of 2, say 7 (0111), n-1 is 6 (0110). bit_width(6) = 3. ceil(log2(7)) = 3.
    // Handles n=1 case correctly: n-1=0, bit_width(0)=0. ceil(log2(1))=0.
    return std::bit_width(n - 1);
}
 
__device__ __forceinline__ size_t device_fast_log2(size_t n)
{
    // For size_t (unsigned long long on many 64-bit systems), digits is 64.
    // __clzll(n) counts leading zeros.
    // If n=8 (0...01000), clzll=60. digits-1-clzll = 64-1-60 = 3. log2(8)=3.
    // If n=7 (0...00111), clzll=61. digits-1-clzll = 64-1-61 = 2. floor(log2(7))=2.
    return std::numeric_limits<size_t>::digits - 1 - __clzll(n);
}
 
__device__ __forceinline__ size_t device_fast_log2_ceil(size_t n)
{
    // If n=8 (0...01000), n-1=7 (0...0111). clzll(7)=61. digits-clzll(n-1) = 64-61 = 3. ceil(log2(8))=3.
    // If n=7 (0...00111), n-1=6 (0...0110). clzll(6)=61. digits-clzll(n-1) = 64-61 = 3. ceil(log2(7))=3.
    // Handles n=1: n-1=0. clzll(0) is undefined but often returns digits (64). digits-digits = 0. ceil(log2(1))=0.
    return std::numeric_limits<size_t>::digits - __clzll(n - 1);
}
 
class CudaTimer
{
public:
    CudaTimer()
    {
        cudaEventCreate(&start);
        cudaEventCreate(&stop);
    }
 
    ~CudaTimer()
    {
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
    }
 
    void start_timer()
    {
        cudaEventRecord(start);
    }
 
    float stop_timer()
    {
        cudaEventRecord(stop);
        cudaEventSynchronize(stop); // Wait for the stop event to complete
        cudaEventElapsedTime(&time, start, stop); // Calculate time difference
        return time;
    }
 
private:
    cudaEvent_t start, stop; 
    float time;              
};
 
template <typename T>
inline T* alloc_to_device(size_t size)
{
    T* d_arr;
 
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMalloc(&d_arr, size * sizeof(T)));
    }
    else
    {
        cudaMalloc(&d_arr, size * sizeof(T));
    }
 
    return d_arr;
}
 
template <typename T>
inline T* copy_to_device(T* d_arr, const T* arr, size_t size)
{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMemcpy(d_arr, arr, size * sizeof(T), cudaMemcpyHostToDevice));
    }
    else
    {
        cudaMemcpy(d_arr, arr, size * sizeof(T), cudaMemcpyHostToDevice);
    }
 
    return d_arr;
}
 
template <typename T>
inline T* copy_to_device(const T* arr, size_t size)
{
    T* d_arr = alloc_to_device<T>(size);
    copy_to_device(d_arr, arr, size);
    return d_arr;
}
 
template <typename T>
inline T* memset_within_device(T* d_arr, size_t size, int value)
{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMemset(d_arr, value, size * sizeof(T)));
    }
    else
    {
        cudaMemset(d_arr, value, size * sizeof(T));
    }
 
    return d_arr;
}
 
template <typename T>
inline T* memset_to_device(size_t size, int value)
{
    T* d_arr = alloc_to_device<T>(size);
    memset_within_device(d_arr, size, value);
    return d_arr;
}
 
template <typename T>
inline T* copy_within_device(T* d_dest, const T* d_src, size_t size)
{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMemcpy(d_dest, d_src, size * sizeof(T), cudaMemcpyDeviceToDevice));
    }
    else
    {
        cudaMemcpy(d_dest, d_src, size * sizeof(T), cudaMemcpyDeviceToDevice);
    }
 
    return d_dest;
}
 
inline void synchronize_device()
{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaDeviceSynchronize());
    }
    else
    {
        cudaDeviceSynchronize();
    }
}
 
template <typename T>
inline T* copy_from_device(const T* d_arr, size_t size)
{
    T* arr = new T[size]; // Allocate host memory
 
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMemcpy(arr, d_arr, size * sizeof(T), cudaMemcpyDeviceToHost));
    }
    else
    {
        cudaMemcpy(arr, d_arr, size * sizeof(T), cudaMemcpyDeviceToHost);
    }
 
    return arr;
}
 
template <typename T>
inline void device_free(T* d_arr)
{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaFree(d_arr));
    }
    else
    {
        cudaFree(d_arr);
    }
}
 
__device__ __forceinline__ size_t get_tid()
{
    return threadIdx.x + blockIdx.x * blockDim.x;
}
 
// Very marginally slower (about 0.5% than get_num_threads() below)
// __device__ __forceinline__ size_t get_num_threads() {
//   return blockDim.x * gridDim.x;
// }
 
 
constexpr size_t get_num_threads()
{
    return BLOCKS * THREADS;
}