mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-22 06:50:25 +08:00
bc310e08d8
* add amd gfx950 mxfp4 support * update the codes for format checking * [AMD][CDNA4] Address review comments for MXFP4 gfx950 support Four issues from CodeRabbit review fixed: 1. hip_fp4.h: Fix FP4 E2M1 denormal decoder round-trip bug. The decoder returned 0.25f for mant==1 (exp==0) but the encoder maps nibble 1 to 0.5f, breaking float→fp4→float round-trips. Fix: return 0.5f to match the encoder and LUT tables. 2. codegen_hip.cc: Fix UB address-of-prvalue in FP4→float16 cast. The generated code took &(__tl_cvt_fp4x2_to_half2(...)), which takes the address of a temporary (C++ UB). Fix: materialize the uint1 return value into a local variable first, matching the pattern already used in the float32/bfloat16 cast paths. 3. codegen_hip.cc: Extend fp4_pair_cast to cover FP4x16 and FP4x32. Previously only 2/4/8-lane FP4 casts were handled; x16/x32 fell through silently to generic vector code that cannot index fp4_e2_* aggregates. The pairwise uint8_t byte logic is correct for any even lane count, so extend to include 16 and 32 lanes. 4. mxfp.py: Narrow bare except to expected import-time errors. The catch-all except Exception silently swallowed all errors from target_is_gfx950(), causing HIP targets to fall back to CUDA/PTX. Fix: catch only ImportError/ModuleNotFoundError/AttributeError. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * update for format checking --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>