CODEWITHSUNDEEP
cuda
rych
rych

Reputation: 712

__device__ __constant__ const

Is there any difference and what is the best way to define device constants in a CUDA program? In the C++, host/device program if I want to define constants to be in device constant memory I can do either

__device__ __constant__ float a = 5;
__constant__ float a = 5;

Question 1. On devices 2.x and CUDA 4, is it the same as,

__device__ const float a = 5;

Question 2. Why is it that in PyCUDA SourceModule("""..."""), which compiles only do device code, even the following works?

const float a = 5;

Upvotes: 4

Views: 11958

Answers (1)

talonmies
talonmies

Reputation: 72372

In CUDA __constant__is a variable type qualifier that indicates the variable being declared is to be stored in device constant memory. Quoting section B 2.2 of the CUDA programming guide

The __constant__ qualifier, optionally used together with __device__, declares a variable that:

In CUDA, constant memory is a dedicated, static, global memory area accessed via a cache (there are a dedicated set of PTX load instructions for its purpose) which are uniform and read-only for all threads in a running kernel. But the contents of constant memory can be modified at runtime through the use of the host side APIs quoted above. This is different from declaring a variable to the compiler using the const declaration, which is adding a read-only characteristic to a variable at the scope of the declaration. The two are not at all the same thing.

Upvotes: 8

Related Questions