Provides utility functions and classes for CUDA operations. More...

#include <bit>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <cuda_runtime.h>
#include <limits>

Include dependency graph for cuda_utils.cuh:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
class	CudaTimer
	A simple CUDA timer class using CUDA events. More...

Macros
#define	BLOCKS 64
	Default number of blocks to launch in CUDA kernels.

#define	THREADS 1024
	Default number of threads per block in CUDA kernels.

#define	gpuErrchk(ans) { gpuAssert(ans, __FILE__, __LINE__); }
	Macro to check CUDA API calls for errors.

Functions
void	gpuAssert (cudaError_t code, const char *file, int line, bool abort=true)
	Asserts if a CUDA error code indicates failure.

__host__ size_t	host_fast_log2 (size_t n)
	Computes the floor of log base 2 for a host integer using std::bit_width.

__host__ size_t	host_fast_log2_ceil (size_t n)
	Computes the ceiling of log base 2 for a host integer using std::bit_width.

__device__ __forceinline__ size_t	device_fast_log2 (size_t n)
	Computes the floor of log base 2 for a device integer using count leading zeros.

__device__ __forceinline__ size_t	device_fast_log2_ceil (size_t n)
	Computes the ceiling of log base 2 for a device integer using count leading zeros.

template<typename T >
T *	alloc_to_device (size_t size)
	Allocates memory on the CUDA device.

template<typename T >
T *	copy_to_device (T d_arr, const T arr, size_t size)
	Copies data from host memory to existing device memory.

template<typename T >
T *	copy_to_device (const T *arr, size_t size)
	Allocates device memory and copies data from host memory to it.

template<typename T >
T *	memset_within_device (T *d_arr, size_t size, int value)
	Sets a region of device memory to a specific byte value.

template<typename T >
T *	memset_to_device (size_t size, int value)
	Allocates device memory and sets it to a specific byte value.

template<typename T >
T *	copy_within_device (T d_dest, const T d_src, size_t size)
	Copies data from one device memory location to another.

void	synchronize_device ()
	Synchronizes the host thread with the default CUDA device stream.

template<typename T >
T *	copy_from_device (const T *d_arr, size_t size)
	Allocates host memory and copies data from device memory to it.

template<typename T >
void	device_free (T *d_arr)
	Frees memory on the CUDA device.

__device__ __forceinline__ size_t	get_tid ()
	Gets the global thread ID within a CUDA kernel launch.

constexpr size_t	get_num_threads ()
	Gets the total number of threads in the kernel launch (compile-time).

Variables
static constexpr bool	CHECK_CUDA_ERRORS = false
	Compile-time flag to enable/disable CUDA error checking via gpuErrchk/gpuAssert.

Detailed Description

Provides utility functions and classes for CUDA operations.

This file includes common CUDA helper functions, error checking macros, memory management wrappers, a timer class, and mathematical utilities for both host and device code.

Definition in file cuda_utils.cuh.

Macro Definition Documentation

◆ BLOCKS

#define BLOCKS 64

Default number of blocks to launch in CUDA kernels.

Definition at line 21 of file cuda_utils.cuh.

◆ gpuErrchk

#define gpuErrchk ( ans ) { gpuAssert(ans, __FILE__, __LINE__); }

Macro to check CUDA API calls for errors.

Calls gpuAssert to check the result of a CUDA API call.

Parameters

ans	The result of the CUDA API call (cudaError_t).

Definition at line 35 of file cuda_utils.cuh.

◆ THREADS

#define THREADS 1024

Default number of threads per block in CUDA kernels.

Definition at line 26 of file cuda_utils.cuh.

Function Documentation

◆ alloc_to_device()

template<typename T >

T * alloc_to_device ( size_t size )

inline

Allocates memory on the CUDA device.

Wrapper around cudaMalloc with optional error checking.

Template Parameters

T	The data type of the array elements.

Parameters

size	The number of elements of type T to allocate.

Returns: T* Pointer to the allocated device memory.

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 190 of file cuda_utils.cuh.

{
    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;
}

◆ copy_from_device()

template<typename T >

T * copy_from_device	(	const T *	d_arr,
		size_t	size
	)

inline

Allocates host memory and copies data from device memory to it.

Wrapper around cudaMemcpy (DeviceToHost) with optional error checking. Allocates host memory using new[].

Template Parameters

T	The data type of the array elements.

Parameters

d_arr	Pointer to the source device memory.
size	The number of elements of type T to copy.

Returns: T* Pointer to the newly allocated host memory containing the copied data.

Note: The caller is responsible for freeing the returned host memory using delete[].; Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 359 of file cuda_utils.cuh.

{
    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;
}

◆ copy_to_device() [1/2]

template<typename T >

T * copy_to_device	(	const T *	arr,
		size_t	size
	)

inline

Allocates device memory and copies data from host memory to it.

Combines allocation (alloc_to_device) and copy (copy_to_device).

Template Parameters

T	The data type of the array elements.

Parameters

arr	Pointer to the source host memory.
size	The number of elements of type T to allocate and copy.

Returns: T* Pointer to the newly allocated and populated device memory.

Note: Uses gpuErrchk indirectly if CHECK_CUDA_ERRORS is true.

Definition at line 245 of file cuda_utils.cuh.

{
    T* d_arr = alloc_to_device<T>(size);
    copy_to_device(d_arr, arr, size);
    return d_arr;
}

◆ copy_to_device() [2/2]

template<typename T >

T * copy_to_device	(	T *	d_arr,
		const T *	arr,
		size_t	size
	)

inline

Copies data from host memory to existing device memory.

Wrapper around cudaMemcpy (HostToDevice) with optional error checking.

Template Parameters

T	The data type of the array elements.

Parameters

d_arr	Pointer to the destination device memory (must be pre-allocated).
arr	Pointer to the source host memory.
size	The number of elements of type T to copy.

Returns: T* The pointer to the destination device memory (d_arr).

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 219 of file cuda_utils.cuh.

{
    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;
}

◆ copy_within_device()

template<typename T >

T * copy_within_device	(	T *	d_dest,
		const T *	d_src,
		size_t	size
	)

inline

Copies data from one device memory location to another.

Wrapper around cudaMemcpy (DeviceToDevice) with optional error checking.

Template Parameters

T	The data type of the array elements.

Parameters

d_dest	Pointer to the destination device memory.
d_src	Pointer to the source device memory.
size	The number of elements of type T to copy.

Returns: T* The pointer to the destination device memory (d_dest).

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 311 of file cuda_utils.cuh.

{
    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;
}

◆ device_fast_log2()

__device__ __forceinline__ size_t device_fast_log2 ( size_t n )

Computes the floor of log base 2 for a device integer using count leading zeros.

Parameters

n	The input value. Must be greater than 0.

Returns: size_t The largest integer k such that 2^k <= n.

Note: Device code only. Uses __clzll (count leading zeros for long long).

Definition at line 103 of file cuda_utils.cuh.

{
    // 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_fast_log2_ceil()

__device__ __forceinline__ size_t device_fast_log2_ceil ( size_t n )

Computes the ceiling of log base 2 for a device integer using count leading zeros.

Parameters

n	The input value. Must be greater than 0.

Returns: size_t The smallest integer k such that 2^k >= n.

Note: Device code only. Uses __clzll (count leading zeros for long long).

Definition at line 118 of file cuda_utils.cuh.

{
    // 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);
}

◆ device_free()

template<typename T >

void device_free ( T * d_arr )

inline

Frees memory on the CUDA device.

Wrapper around cudaFree with optional error checking.

Template Parameters

T	The data type of the array elements (used for pointer type).

Parameters

d_arr Pointer to the device memory to free.

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 385 of file cuda_utils.cuh.

{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaFree(d_arr));
    }
    else
    {
        cudaFree(d_arr);
    }
}

◆ get_num_threads()

constexpr size_t get_num_threads ( )

constexpr

Gets the total number of threads in the kernel launch (compile-time).

Calculates the total number of threads based on the predefined BLOCKS and THREADS macros.

Returns: constexpr size_t The total number of threads.

Note: This is a constexpr function, evaluated at compile time.

Definition at line 424 of file cuda_utils.cuh.

{
    return BLOCKS * THREADS;
}

◆ get_tid()

__device__ __forceinline__ size_t get_tid ( )

Gets the global thread ID within a CUDA kernel launch.

Calculates the unique ID for the current thread across all blocks in the grid. Assumes a 1D grid and 1D block configuration.

Returns: size_t The global thread identifier.

Note: Device code only.

Definition at line 406 of file cuda_utils.cuh.

{
    return threadIdx.x + blockIdx.x * blockDim.x;
}

◆ gpuAssert()

void gpuAssert	(	cudaError_t	code,
		const char *	file,
		int	line,
		bool	abort = `true`
	)

inline

Asserts if a CUDA error code indicates failure.

Checks the provided CUDA error code. If it's not cudaSuccess, it prints an error message including the error string, file name, and line number. If abort is true, it terminates the program.

Parameters

code	The CUDA error code to check.
file	The source file where the check occurs (FILE).
line	The line number where the check occurs (LINE).
abort	If true (default), exits the program upon error.

Definition at line 49 of file cuda_utils.cuh.

{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
 
        if (abort)
        {
            exit(code);
        }
    }
}

◆ host_fast_log2()

__host__ size_t host_fast_log2 ( size_t n )

inline

Computes the floor of log base 2 for a host integer using std::bit_width.

Parameters

n	The input value. Must be greater than 0.

Returns: size_t The largest integer k such that 2^k <= n.

Note: Host code only.

Definition at line 74 of file cuda_utils.cuh.

{
    // 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_fast_log2_ceil()

__host__ size_t host_fast_log2_ceil ( size_t n )

inline

Computes the ceiling of log base 2 for a host integer using std::bit_width.

Parameters

n	The input value. Must be greater than 0.

Returns: size_t The smallest integer k such that 2^k >= n.

Note: Host code only.

Definition at line 88 of file cuda_utils.cuh.

{
    // 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);
}

◆ memset_to_device()

template<typename T >

T * memset_to_device	(	size_t	size,
		int	value
	)

inline

Allocates device memory and sets it to a specific byte value.

Combines allocation (alloc_to_device) and memory set (memset_within_device).

Template Parameters

T	The data type of the array elements.

Parameters

size	The number of elements of type T to allocate and set.
value	The integer byte value to set (e.g., 0).

Returns: T* Pointer to the newly allocated and initialized device memory.

Note: Uses gpuErrchk indirectly if CHECK_CUDA_ERRORS is true.

Definition at line 291 of file cuda_utils.cuh.

{
    T* d_arr = alloc_to_device<T>(size);
    memset_within_device(d_arr, size, value);
    return d_arr;
}

◆ memset_within_device()

template<typename T >

T * memset_within_device	(	T *	d_arr,
		size_t	size,
		int	value
	)

inline

Sets a region of device memory to a specific byte value.

Wrapper around cudaMemset with optional error checking.

Template Parameters

T	The data type of the array elements (used for size calculation).

Parameters

d_arr	Pointer to the device memory region to set.
size	The number of elements of type T in the region.
value	The integer byte value to set (e.g., 0).

Returns: T* The pointer to the device memory region (d_arr).

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 265 of file cuda_utils.cuh.

{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaMemset(d_arr, value, size * sizeof(T)));
    }
    else
    {
        cudaMemset(d_arr, value, size * sizeof(T));
    }
 
    return d_arr;
}

◆ synchronize_device()

void synchronize_device ( )

inline

Synchronizes the host thread with the default CUDA device stream.

Waits until all preceding commands in the default stream have completed. Wrapper around cudaDeviceSynchronize with optional error checking.

Note: Uses gpuErrchk if CHECK_CUDA_ERRORS is true.

Definition at line 333 of file cuda_utils.cuh.

{
    if constexpr (CHECK_CUDA_ERRORS)
    {
        gpuErrchk(cudaDeviceSynchronize());
    }
    else
    {
        cudaDeviceSynchronize();
    }
}

Variable Documentation

◆ CHECK_CUDA_ERRORS

constexpr bool CHECK_CUDA_ERRORS = false

staticconstexpr

Compile-time flag to enable/disable CUDA error checking via gpuErrchk/gpuAssert.

Note: Set to true to enable runtime checks (potentially slower), false to disable them.

Definition at line 66 of file cuda_utils.cuh.

Classes

Macros

Functions

Variables

Detailed Description

Macro Definition Documentation

◆ BLOCKS

◆ gpuErrchk

◆ THREADS

Function Documentation

◆ alloc_to_device()

◆ copy_from_device()

◆ copy_to_device() [1/2]

◆ copy_to_device() [2/2]

◆ copy_within_device()

◆ device_fast_log2()

◆ device_fast_log2_ceil()

◆ device_free()

◆ get_num_threads()

◆ get_tid()

◆ gpuAssert()

◆ host_fast_log2()

◆ host_fast_log2_ceil()

◆ memset_to_device()

◆ memset_within_device()

◆ synchronize_device()

Variable Documentation

◆ CHECK_CUDA_ERRORS