Go to the source code of this file.
Macros | |
| #define | NEKO_WAVE_SIZE 64 /* host pass, see the note above */ |
| #define | NEKO_WAVE_SIZE_UNIFORM 64 |
Macro Definition Documentation
◆ NEKO_WAVE_SIZE
Wavefront width (HIP)
GCN and CDNA run 64 lane wavefronts, RDNA runs 32, and gfx10 and later can be built either way with -mwavefrontsize32/-mwavefrontsize64. The warp level shuffle reductions have to know the width at compile time, so that the ladder stays the same fully unrolled instruction sequence it has always been on CDNA rather than becoming a runtime loop.
warpSize is not usable for that: in device code it expands to __builtin_amdgcn_wavefrontsize(), which is only folded to a literal in the backend, long after the preprocessor and the template instantiation that would need it.
IMPORTANT: NEKO_WAVE_SIZE is resolved per compilation pass, and hipcc compiles every .hip twice. None of the macros it keys off exist in the host pass, so the host falls back to 64 whatever the device is. It may therefore only be used inside device code. Anything that has to agree between the two passes – launch geometry, kernel template arguments, launch_bounds – must use NEKO_WAVE_SIZE_UNIFORM instead, or the host will launch a kernel instantiation the device pass never emitted.
