News from GPU parameter tuning

Investigated on the optimal configurations found by the tuner.

launch bounds "issue"

HIP definition of launch_bounds (link): 

__global__ void __launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_WARPS_PER_EXECUTION_UNIT)
MyKernel(hipGridLaunch lp, ...)
...

MIN_WARPS_PER_EXECUTION_UNIT - directs the compiler to minimize resource usage so that the requested number of warps can be simultaneously active on a multi-processor.

Compute Unit (CU) for AMD = Streaming Multiprocessors for Nvidia. A CPU is composed of one or more Execution Units (EU) which are responsible for executing waves. For MI50: 60 CUs, 4 EUs per CU.

CUDA definition of __launch_bounds:

__launch_bounds(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MULTIPROCESSOR)

"The second parameter __launch_bounds parameters must be converted to the format used __hip_launch_bounds, which uses warps and execution-units rather than blocks and multi-processors"

Implication

This means that when the launch bounds are set on the parameter file and compiled for AMD GPUs, we set the min number of active warps each EU has to contain

Example: with #define GPUCA_LB_GPUTPCGMMergerCollect 512, 9

SGPRs: 62
VGPRs: 28
ScratchSize [bytes/lane]: 22464
Dynamic Stack: False
Occupancy [waves/EU]: 9
SGPRs Spill: 28
VGPRs Spill: 111
LDS Size [bytes/block]: 0

So if GetGridAuto is used to launch the kernel, the grid will be MIN_WARPS_PER_EXECUTION_UNIT * CUs, and the compiler will reserve registers accordingly to the specified MIN_WARPS_PER_EXECUTION_UNIT

On MI50, hw limit is 10 active warps per EU --> when the second parameter is > 10 in the launch bounds, it is ignored and the compiler discards this hint, usually trying not to spill any register

Example: with #define GPUCA_LB_GPUTPCGMMergerCollect 512, 15

SGPRs: 104
VGPRs: 73
ScratchSize [bytes/lane]: 22272
Dynamic Stack: False
Occupancy [waves/SIMD]: 3
SGPRs Spill: 0
VGPRs Spill: 0
 LDS Size [bytes/block]: 0

Each thread uses 73 VGPRs, total of 256 VGPRs per thread --> can accomodate 3 warps without spilling

Grid size = 600 mistery solved

This is why grid size 600 is such a magic number in the tuning, and why grid size is so important. In reality, it is steering the number of active warps per EU. After 10 active warps (600/CUs), the compiler discard the hint, thus changing optimization strategy.

Sometimes spilling and increasing occupancy helps, sometimes it is better to reduce the spilling.

What to do with this