Built-in Variables Memory Location ? in which memory are built in variables stored

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

In which memory space are Built-in Variables stored?

in global, local, shared, registers ?

such as: gridDim, blockIdx, blockDim, threadIdx, warpSize.

So if I access threadIdx.x or blockDim.x for example very often during a kernel, should I better copy it to a register?

//use threadIdx.x directly multiple times

vs

unsigned int threadId = threadIdx.x;

//use threadIdx multiple times

There are no informations about this in the Nvidia programming guide!

Till now I only used new registers if my desired threadId needs to be calculated and cannot directly retreived by for example threadIdx.x

Edit: In some other threads I found they should reside in shared memory, so this thread can be closed

2

the programming guide says, they are stored in special registers.

but i dont think there will be much difference between using threadIdx.x everywhere or a dedicated int-variable, the compiler optimizes it the right way

3

Two minor notes, just to clear up some misconceptions:

  • Many architectures, include CUDA, can only operate on things in registers. Regardless of where threadIdx.x is stored, it will need to be copied to a register to be used in a calculation or as an array index. How long it stays in a register is up to the compiler.

  • Relating to the previous point, there is no one-to-one relationship between C variables and registers. The compiler (and PTX assembler in CUDA) decides what to store in a register and for how long. A single expression might require several registers to compute, but after the value is used, the compiler is free to throw it away and use the register for something else. If there are not enough registers to hold the number of simultaneously required values, then the compiler will push some of the values into “local memory”, which is just a section of global memory where storage can be set aside for each thread.

4

Thanks for your answer, cleared some things up for me.