mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Zhang JasonandLeiWang1999
3d95d65000
* [AMD][RDNA4] Fix RDNA4 (gfx1201 / Wave32) CI failures
- reduce.h: Replace hardcoded wave64 assumptions with runtime
__builtin_amdgcn_wavefrontsize() in SharedReduceWarp and AllReduce
- scan.h: Add SEG=32 support in InclusiveScan/CumSum/CumMax for RDNA wave32
- finalize_reducer.cc: Use TargetGetWarpSize() instead of hardcoded 64
- carver/arch/__init__.py + rdna.py: Add gfx12 (RDNA4) support alongside gfx11
- wmma_macro_generator.py: Fix k_pack indexing in ldmatrix_a/b for k_pack > 1
These changes correspond to PR #2210 (Fix RDNA4 gfx1201/Wave32 CI Failures),
rebased cleanly onto upstream main (a8d93798).
* [AMD][HIP] Fix HIPMath dispatch for vector dtype in T.Parallel
When BLOCK_M > threads, TileLang generates vector types (e.g. float32x4)
for T.Parallel loops. HIPMath::operator() only handled scalar types, so
vector dtypes fell through to return causing Unresolved call errors
in codegen for math intrinsics like tirx.exp2.
Fix: add a vector-type early-return that recurses on element_of(), so
float32xN correctly resolves to exp2f the same as scalar float32.
This complements the tirx.* registration fixes in #2282 and enables
Online Softmax on gfx1151 with BM > threads (e.g. BM=1024, TH=256).
* [Review] Fix 3 issues found in code review
1. tilelang/carver/arch/__init__.py: error message was not updated
when gfx12 support was added — still said gfx11 only. Now says
gfx11/gfx12 targets only, consistent with rdna.py.
2. src/tl_templates/hip/reduce.h: the static_assert(Threads % kWarpSize == 0)
was removed when kWarpSize became a runtime value, leaving no safety
check. Added equivalent runtime assert().
3. src/tl_templates/hip/scan.h: run_auto() else branch silently fell
through to SEG=64 for any non-32 wavefront size. Added assert to
catch unsupported wavefront sizes early.
* update codes for format checking
* [Review] Address review feedback and rebase on upstream main
- intrin_rule_hip.cc: remove redundant HIPMath scalar recursion; upstream
PR #2286 already handles vector dtypes via DispatchPureExternScalarized
at the IR level (correct for HIP packed carrier types)
- wmma_macro_generator.py: rename l/r to row_off/col_off to fix Ruff E741
ambiguous variable name warnings
- scan.h: guard run_auto() branches with threads>=32 / threads>=64 so a
threads==32 kernel on Wave64 doesn't compute TILE=0 and divide by zero
- reduce.h: replace assert (stripped in Release) with unconditional early
return when Threads % kWarpSize != 0, preventing silent num_warps==0
* update codes for format checking
* fix
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>