How does dynamic reassignment of register capacity among warp-groups on Hopper cards work?

[mrakgr]

In CUDA C++ Programming Guide the guide mentions dynamic reassignment of register capacity among warp-groups to support even larger matrices. I didn’t see any special instructions for this in the PTX guide and am looking for more info on how this works. I got access to a HGX machine and will be optimizing ML models on it in the future, so this topic is something I need to familiarize myself with.

Any help would be appreciated.

[striker159]

You would use the PTX instruction setmaxnreg: PTX ISA 8.5

See the explanation by Greg in this post:

Sm90: setmaxnreg will change Occupancy dynamically?

A CUDA grid launch specifies the number of registers per warp. All warps in the CTA (thread block) launch with the same number of registers. The PTX instruction setmaxnreg allows a warp (e.g. producer warp) to return registers to the CTA register pool (setmaxnreg.dec) and other warps to acquire registers from the CTA register pool (setmaxnreg.inc). The total number of registers allocated to the CTA does not change. The total number of registers allocated to all warps in the CTA cannot exceed th…