register spilling

[ee06b109]

I have a kernel and launches with 64 threads and 28 blocks.
suppose my GPU has x registers in each SM.
Will register spilling occur when one of the thread consume more than x/64 registers?
Thanks

Example:
if(threadIdx.x == 0)
int i[x/32];

[Easyhard]

Hi,
I belive each SM will launch as much blocks as rescourse limits can be satisfied, so I guess in your situation each SM launch one block one time, and no spilling occurs.
If I make any mistake, PLZ let me know^_^

easyhard

[ee06b109]

Hi,

I belive each SM will launch as much blocks as rescourse limits can be satisfied, so I guess in your situation each SM launch one block one time, and no spilling occurs.

If I make any mistake, PLZ let me know^_^

easyhard

Someone said that if i launch 64 threads in a block, the registers will be divided into 64 sets. Then, if one thread consumes more than x/64 registers, the spilling will occur.

I’m not sure whether it’s true.

Thank you for your reply~

[tera]

This is only an upper bound for the available number of registers. You need to consider all threads of all blocks currently active on an SM, and take into account that registers are allocated in chunks of more than one register. Check section 4.2 of the Programming Guide for the details.