mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-22 06:50:25 +08:00
Chennes
0f13a2c548
CUDA lowering emits hpow and hfmod for 16-bit pow and fmod, but no CUDA toolkit declares either name, so fp16/bf16 fmod and variable-exponent fp16 pow fail to compile. Integer-exponent pow is unaffected: TileLang expands it into repeated multiplication before lowering. Add half_t and bfloat16_t overloads that evaluate with powf or fmodf in float32 and convert back, matching the existing bridge pattern. Remove math.h's hpow-to-powf macro: it shadows the new overload at the call site, and compiles only where the result is assigned directly. Add codegen and numerical coverage for fp16 pow and both fmod dtypes. The bfloat16 hpow bridge completes the backend layer, while T.pow on bfloat16 remains blocked by an upstream TVM dtype check.