tilelang

mirror of https://github.com/tile-ai/tilelang.git synced 2026-10-23 07:00:29 +08:00

Files

T

Zhang Jason 53a4c9866b [AMD] [gfx950]Fix multiple HIP codegen bugs to support TileKernel (#2099 )

* Fix HIP codegen for sync_warp, sync_grid, and local.var initialisation

* [AMD/HIP] Fix warp_reduce VGPR bug, ShuffleNode packing, and Pipelined LDS overflow

Extends PR #2096 with three additional fixes for CDNA (MI350) targets:

Fix 1 — src/tl_templates/hip/reduce.h: warp_reduce width=32
  The old 6-step butterfly called __shfl_xor(value, 32) without a width
  argument. On CDNA (wave64) with 32 active threads, lanes 32-63 are
  inactive and hold uninitialised VGPRs, producing NaN in reduce_max /
  reduce_sum / AllReduce. Fix: remove the step-32 shuffle; pass width=32
  to all remaining 5 steps so every shuffle stays within the [0,31] group.

Fix 2 — src/target/codegen_hip.cc + src/tl_templates/hip/common.h:
         ShuffleNode bfloat16x2 / float16x2 packing
  CodeGenC emitted `uint1(a, b)` for bfloat16x2 construction, which is an
  invalid HIP constructor call. Fix: override VisitExpr_(ShuffleNode) in
  CodeGenTileLangHIP to emit `uint1{__pack_bfloat162(a, b)}` / `uint1{
  __pack_half2(a, b)}` using aggregate initialisation. Also add five
  bfloat16x2 math overloads for uint1 carrier (abs2/max2/min2/add2/mul2).

Fix 3 — src/transform/pipeline_planning.cc: skip T.Pipelined(num_stages>1)
  Double-buffering doubled LDS per loop-body buffer. On CDNA (≤128 KB LDS
  per workgroup), this caused hipModuleLaunchKernel EINVAL. Fix: when
  TargetIsRocm() && num_stages > 1, skip pipeline planning and fall back
  to a plain sequential loop with synchronous T.copy.

Also: fix __habs(hip_bfloat16) and __habs(float16_t) in common.h to use
__builtin_memcpy instead of reinterpret_cast to avoid strict-aliasing UB
(as flagged by CodeRabbit on PR #2096).

Tests: 19 new cases added to testing/python/amd/test_tilelang_hip_codegen.py
covering all three fixes. All 42 tests pass on MI350 (gfx950).

* [AMD/HIP] Merge test_tilelang_hip_bugfixes.py into test_tilelang_hip_codegen.py

Consolidate all HIP regression tests into a single file.  The merged file
covers all six fixes with 32 tests total (previously split across two files
with duplicated test cases for warp_reduce, pipelined GEMM, and ShuffleNode).

Changes versus the two individual files:
- Deduplicated test_warp_reduce_no_nan (identical in both files)
- Deduplicated test_pipelined_no_lds_overflow / test_pipelined_shared_mem_no_launch_error
- Deduplicated test_pipelined_multi_stage_fp16_gemm
- Merged bfloat16 shuffle tests: source check + runtime correctness in one function
- Kept PR #2096 source-inspection tests (alloc_var, sync_warp, sync_grid)
- Added runtime tests from bugfixes: inf init, serial loop accumulation,
  float scalar readback, two-group wave64 reduce, float16 shuffle

* fixup: correct LDS size comment — gfx950 has 160 KB, not 128 KB

gfx942 (CDNA3 / MI300X) has 64 KB LDS per workgroup.
gfx950 (CDNA4 / MI350)  has 160 KB LDS per workgroup (see PR #2058).

The old comment said '128 KB' which is wrong for both generations.
Updated pipeline_planning.cc and the test docstrings to reflect the
correct per-architecture limits.

* update for format checking

2026-04-27 02:31:01 +08:00

amd

[AMD] [gfx950]Fix multiple HIP codegen bugs to support TileKernel (#2099 )

2026-04-27 02:31:01 +08:00

analysis

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002 )

2026-04-07 23:54:55 +08:00

arith

[Refactor] Unify @jit and @lazy_jit into a single @jit decorator (#1632 )

2026-01-09 15:41:39 +08:00

autotune

[Cache] Refactor cache namespace layout (#2057 )