mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Chennes
e54cbabb12
T.shfl_sync, T.shfl_xor, T.shfl_down, and T.shfl_up lower to CUDA's raw __shfl_*_sync builtins. CUDA has no exact overloads for the narrow floating-point wrappers TileLang emits, so calls either resolved to the float overload and returned a float that could not copy-initialize the wrapper, or were ambiguous. Either way the shuffle failed to compile. Add exact non-template overloads for half_t, bfloat16_t, float_e4m3_t, and float_e5m2_t. The overloads carry the raw bits through CUDA's uint32_t shuffle and preserve mask, lane, delta, and width semantics. A local macro generates the four builtin overloads per type and is undefined immediately after use, leaving codegen and native CUDA overload resolution unchanged. Add exact lane-permutation coverage for all four shuffle operations and all four wrapper types in both materialized-temporary and inline forms. Fixes #2569