Using constants with CUDA

Question

Which is the best way of using constants in CUDA?

One way is to define constants in constant memory, like:

// CUDA global constants
__constant__ int M;

int main(void)
{
    ...
    cudaMemcpyToSymbol("M", &M, sizeof(M));
    ...
}

An alterative way would be to use the C preprocessor:

#define M = ... 

I would think defining constants with the C preprocessor is much faster. Which are then the benefits of using the constant memory on a CUDA device?

Answer 1

1. constants that are known at compile time should be defined using preprocessor macros (e.g. #define) or via C/C++ const variables at global/file scope.
2. Usage of __constant__ memory may be beneficial for programs who use certain values that don't change for the duration of the kernel and for which certain access patterns are present (e.g. all threads access the same value at the same time). This is not better or faster than constants that satisfy the requirements of item 1 above.
3. If the number of choices to be made by a program are relatively small in number, and these choices affect kernel execution, one possible approach for additional compile-time optimization would be to use templated code/kernels