[Python] CUDA Registers and Offloading to Shared Memory [closed]

Edit update: With the help of a commenter below, I uncovered two issues:

1. A function was marked virtual for a base cpp class I was using, in particular the ArrayMetrics.Size() and ArrayMetrics.push_back() per the messages below. There are several topics on challenges with virtual functions in __device__code (How to implement device side CUDA virtual functions?)
2. I still had 232 bytes stack frame being allocated w/ some gmem remaining. Once I unrolled two small loops, I got the following:

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 48 registers, used 0 barriers, 480 bytes cmem[0]


Original Question:

I have been working on a cuda script where after profiling with Nsight, I can see where in the source it's currently being allocated to global memory. I am currently running this on a 4090 GPU.

I turned on the compile flags to also tell me the local memory usage and it is as follows:

/home/mem/scripts/run-cuda.py:253: UserWarning: The CUDA compiler succeeded, but said the following:
ptxas info : Overriding global maxrregcount 127 with entry-specific value 128 computed using thread count
ptxas info : 40 bytes gmem, 8 bytes cmem[4]
ptxas info : Compiling entry function 'cuda_entry' for 'sm_89'
ptxas info : Function properties for cuda_entry
232 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 52 registers, used 0 barriers, 232 bytes cumulative stack size, 480 bytes cmem[0]
ptxas info : Compile time = 31.596 ms
ptxas info : Function properties for _ZN17ArrayMetrics12valid_metEv
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for _ZN17ArrayMetrics4sizeEv
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for _ZN17ArrayMetrics9push_backEO5Met
16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads


So if I am understanding this correctly:

1. 40 bytes gmem, 8 bytes cmem[4] I have 40 bytes ending up in global memory and 8 bytes ending in constant memory.
2. The "cuda_entry" function uses 232 bytes in its stack frame, but nothing is spilled over to global memory?
3. Used 52 registers, used 0 barriers, 232 bytes cumulative stack size, 480 bytes cmem[0] - Only 52 active registers
4. In the ArrayMetrics.push_back(Met) line listed, that seems unable to fit in registers, so we have the local memory spilling?

To try to minimize the spilling from registers into local mem, I have added the following, but it does not seem to change the actually used 52 registers:

1. Compilation flag: '--maxrregcount=127' (I put 127 since I wanted to see it forcing the 128 as pictured above).
2. I have also tried the directive: __launch_bounds__(512, 1) to try to tell the compiler that the # of threads will be limited to 512 so that hopefully more registers per SM will be allocated to each thread

I understand increasing register allocation will decrease active threads/blocks but wanted to see if keeping the entire computation in registers would improve performance.

Alternatively, I am debating if I allocated a shared memory array to house each of the ArrayMetrics instances for each thread would possibly help performance?

Resources consulted:

• Limiting register usage in CUDA: __launch_bounds__ vs maxrregcount
• Counting registers/thread in Cuda kernel

Continue reading...