* [BugFix] Skip DecoupleTypeCast on Evaluate roots to keep cp.async operands in their address spaces
* [Test] Assert cp_async_gs lines exist before checking their operands
The codegen regression test filtered kernel source with
, so a source with no global->shared calls
left the operand assertion with nothing to inspect and the test passed
vacuously. The preceding did not catch
that: cp_async_commit/cp_async_wait come from the pipeline's sync
skeleton, a separate emit site, and match that substring on their own.
Collect the cp_async_gs lines, assert the collection is non-empty, then
check each one for _local_cast operands.
A large number of tests import CUDA-only internals directly (for example
`tilelang.cuda.pipeline.CUDAPassPipelineBodyPrologue`) or exercise
NVIDIA-specific features such as TMA, LDGSTG, mbarrier, warp specialization,
tcgen05 and tfloat32. `src/cuda/CMakeLists.txt` only compiles a small
"ALWAYS" subset when `USE_CUDA=OFF`, so those FFI symbols are never
registered and the tests fail with:
AttributeError: module 'tilelang.cuda._ffi_api' has no attribute
'AnnotateDeviceBoundTmaCopies'
This is not ROCm-specific -- the same tests fail on any build without the
CUDA backend, including Metal and CPU-only.
Add `@tilelang.testing.requires_cuda` to those tests, plus three
portability fixes that keep coverage on non-CUDA backends instead of
skipping it:
* `test_tilelang_transform_layout_inference.py`: CUDA emits
`__launch_bounds__(N, 1)` while HIP emits `__launch_bounds__(N)`.
Match both.
* `test_tilelang_transform_pipeline_planning.py`: PipelinePlanning
deliberately strips pipeline annotations on ROCm targets other than
gfx950 (see the guard in src/transform/pipeline_planning.cc), so assert
the documented sequential fallback there.
* `test_tilelang_kernel_gemm.py`: the two f32f32f32 cases use `T.tfloat32`,
which is NVIDIA-only, so narrow `requires_cuda_or_cdna` to
`requires_cuda`.
Also refresh three stale RDNA tests: gfx12 is now a supported generation
(`get_arch` accepts 11 and 12), so they can no longer assert that gfx1200
is rejected. Force an unsupported generation instead, which is the only
way to reach the guard now that `target_is_rdna` reports only gfx11/gfx12.
Verified on 8x MI300X (gfx942, ROCm 7.2 container, torch nightly rocm7.2):
the ROCm suite goes from 105 failures to 0 with no loss of passing tests.
The VerifyParallelLoop data race check can report false positives on
shared buffer stores whose per-thread addresses cannot be proven
distinct (e.g. indices loaded from fragments). Disable it by default;
users can opt in via the TILELANG_ENABLE_DATA_RACE_CHECK environment
variable or by explicitly setting the tl.disable_data_race_check pass
config to False.
* [CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas
LowerSharedTmem gave every T.alloc_tmem buffer its own tcgen05.alloc,
rounded up to a power of two. Small buffers therefore each burned a
32-column allocation: a 384-column accumulator next to three 4-column
block-scale buffers needs 396 columns but asked the hardware for
512 + 32 + 32 + 32 = 608 of the 512 a CTA has, so that kernel could not
be expressed at all. Kernels worked around it by declaring one wide
buffer and slicing scale factors out of it by hand.
Plan the allocations instead. PlanTmemArenas packs logical buffers into
shared allocations, and each buffer's column offset inside the one it
shares is folded into every address it forms -- both the encoded
coordinate of a BufferLoad, which tcgen05.ld/st consume, and the
(base, offset) argument pairs of the tl.ptx_tcgen05_* intrinsics.
Buffers at least 32 columns wide keep the 32-column alignment
tcgen05.alloc would have given them alone; narrower ones use the
4-column granularity of tcgen05.cp.32x128b.warpx4, the same boundary
CUTLASS and DeepGEMM place these operands on.
A buffer joins an existing allocation only when that strictly reduces the
total column count, so packing can never grow a kernel's TMEM footprint
and a kernel whose buffers already fit lowers exactly as before. The
mxfp8 block-scaled GEMM drops from 320 to 288 columns; flash attention
and single-accumulator GEMMs are untouched.
Buffers with an explicit T.deallocate_tmem keep their own allocation,
since releasing one early would release whatever shares it. A base
address that reaches an operation this pass cannot offset is rejected
rather than silently addressing the start of the allocation. A block
that allocates more than 512 columns is now reported at compile time
instead of failing inside tcgen05.alloc on device.
* [Test] Cover TMEM arena planning
Pins the plan the packer produces, not just that it produces one:
- 396 columns across four buffers become one 512-column allocation at
offsets 0 / 384 / 388 / 392, matching DeepGEMM's hand-written layout
for this kernel column for column
- block-scale factors share an allocation while the accumulator keeps
its own, because packing all three would cost more than it saves
- equal-cost cases keep separate allocations, so a kernel that already
fits is untouched
- allocations are issued widest first, as PTX requires
- packing never allocates more than separate allocations would, over six
shape combinations
- an explicit T.deallocate_tmem keeps its own allocation
- a dynamic pipeline-stage coordinate still carries the column offset
- a block over the 512-column budget is reported
* [Test] Roundtrip two TMEM buffers that share one allocation
Covers the tcgen05.st/ld address path for a buffer packed at a nonzero
column offset, which no existing kernel exercises: the block-scale
buffers that share an allocation today are never read back into
registers.
The wide buffer's roundtrip is the assertion that matters. The narrow
one would still round-trip if its store and its load dropped the column
offset together -- but its store would land on the wide buffer's first
columns, so only the wide readback catches it. Verified by mutation:
dropping the offset in LowerSharedTmem fails this test.
* [Test] Guard the TMEM transform tests on Blackwell
Gate every TMEM test the same way, matching the decorators the tcgen05
language and kernel tests already carry.
These three lower IR for a hardcoded sm_100 target and need no device, so
this does trade running them on any machine for consistency with the rest
of the TMEM suite.
* [BugFix] Skip descriptor TMA for device-bound copy bases
* [BugFix] Run TMA base annotation independently of WS
* [Refactor] Expose TMA base annotation as a standalone pass
* [Docs] Restore warp-specialization pipeline comment
* [Transform] Never classify side-effecting binds as replayable
IsReplayableScalarBind (introduced in #2245) decides whether a scalar
bind may be dissolved and its value re-evaluated at every use site
(IfStmtBinding, default-on) or replayed across pipeline stages
(PipelinePlanning, InjectSoftwarePipeline). The predicate checked the
dtype and read/write disjointness, but never the value's effect kind,
so a bind like
pos = T.atomic_add(counter[0], 1, return_prev=True)
used at two sites (bound check + store index) was inlined into both use
sites, executing the atomic once per use site instead of exactly once
and silently corrupting compaction/histogram/top-k kernels since
v0.1.11. The atomic intrinsic is correctly registered kOpaque; its
pointer argument enters the expression through tvm_access_ptr, which
the access collector records as a read, so the write-disjointness check
passed and the bind was declared replayable.
Reject any bind whose value has SideEffect > kReadState, mirroring the
effect guard StmtSimplifier already applies to let inlining. Read-only
values remain replayable; read consistency is enforced by the existing
write-set disjointness check.
The shipped examples/deepseek_v32/topk_selector.py uses exactly this
pattern and is miscompiled on v0.1.11/v0.1.12; its test only printed
recall and could not fail, so it also gains a recall assertion.
Regression tests:
- transform: IfStmtBinding must keep a bind whose value is an atomic
RMW returning the previous value.
- language (CUDA): compaction kernel using atomic_add(return_prev=True)
must contain exactly one AtomicAddRet and produce a permutation with
counter == N.
* [Transform] Validate scalar bind replay safety
Track access-pointer write masks during bind and pipeline analysis, and reject unsafe precomputed replay annotations during pipeline injection.
Add regressions for atomic target reads and side-effecting replay masks.
* Remove recall calculation from top-k selector
Removed recall calculation and assertion for top-k selection.
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>
* [BugFix] Track GEMM accumulator writes in warp-specialization liveness collector
* [BugFix] Track scale-factor regions and region extents in gemm liveness
* [BugFix] Return previous value for atomic min/max/add (scalar + vectorized)
`atomic_min`/`atomic_max` (scalar) and `atomic_addx2`/`atomic_addx4`
(vectorized) typed their call with a non-void return under
`return_prev=True` but always selected the void `*_elem_op`, so codegen
emitted an empty binding (`prev = ;`) and nvcc failed.
Select the `*_ret_elem_op` variant under `return_prev`, as `atomic_add`
already does:
- min/max: the ret ops, codegen handlers, and device templates
(AtomicMaxRet/AtomicMinRet) already existed.
- addx2/addx4 (float32): register the ret ops (builtin.cc/.h) and add
codegen handlers emitting AtomicAddx2Ret/AtomicAddx4Ret (already in
the device headers) as expressions; the return is typed as the packed
2-/4-lane vector so the store matches.
- addx2 (fp16/bf16): the ret store LHS is `uint1`; codegen wraps the
half2/bf162 result with `tl::to_uint1`.
- addx4 (fp16/bf16): add AtomicAddx4Ret(half_t*)/(bfloat16_t*) device
overloads that run two AtomicAddx2Ret halves and return them packed as
uint2 (a fp16x4 is stored as uint2).
The 16-bit x4 path is per-pair atomic (two half2 atomics), not
whole-quad single-atomic -- hardware has no vector atomic wider than
`.f16x2` for 16-bit types. This matches the existing fp32-x4 scalar
fallback, and the docstring only promises "quad-width atomic addition".
Adds tests across scalar/x2/x4 and float32/fp16/bf16.
Fixes#2667
* minor fix for missing conditiions
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* Fix physical wait counts for async pipeline groups
* Add commit group tracking for async pipeline in inject_pipeline.cc
- Introduced buffer_to_commit_group mapping and commit_group_count in AsyncStateGlobal.
- Enhanced wait handling for tail consumers in async pipeline to ensure correct physical wait counts.
- Updated tests to validate fine-grained physical waits and tail drain waits for async pipeline.
* [Refactor] Remove unsupported tl_gemm and tl_gemm_sp operations
* Removed references to tl_gemm and tl_gemm_sp from various files, including code generation and built-in definitions, as these operations are currently unsupported.
* Updated related comments and diagnostics to reflect the removal of these operations, ensuring clarity in the codebase.
* This change simplifies the handling of unsupported operations and improves overall code maintainability.
* Reduce CUDA template include overhead
* Fix CUDA reduce fast min max helpers
* lint fix
* Fix CUDA swizzle ceil div helper
* Restore CUDA wait_wgmma intrinsic helper
* Fix CUDA cython TMA host adapter include
* Remove unnecessary check for Cutlass fast math header in CUDA code generation
* Update LibraryGenerator to correctly handle CUDA gencode flags for shared-library compilation, ensuring compatibility with target architecture and avoiding Hopper-only instruction issues.
* Refactor smem_ptr_to_uint function to use __cvta_generic_to_shared for improved clarity and performance.
* Refactor example scripts and CUDA math intrinsics for improved performance and clarity
- Updated `example_dequant_gemm_fp4_hopper.py` to directly call `run_regression_perf()` and print latency, removing unused argument parsing code.
- Modified `example_topk.py` to specify the backend as "cupti" in the benchmarking function.
- Enhanced CUDA code generation by introducing `RequiresTileLangMathHeader` to manage math header dependencies for specific functions.
- Updated barrier operations in `barrier.h` to use the correct syntax for shared memory barriers.
- Improved testing for bfloat16 math intrinsics by replacing pytest skips with decorators for CUDA availability checks.
* Refactor CUDA code for improved readability and consistency
- Reformatted the `RequiresTileLangMathHeader` function for better readability.
- Updated assembly syntax in `barrier.h` for clarity and consistency.
- Cleaned up whitespace in the bfloat16 math intrinsics test file to enhance code cleanliness.
* Remove unnecessary whitespace in CUDA code for improved readability
* Add TileLang intrinsic for prefetching TMA descriptor
- Introduced `prefetch_tma_descriptor` intrinsic in C++ and updated relevant files to use `call_intrin` instead of `call_extern`.
- Updated CUDA code generation to handle the new intrinsic correctly.
- Modified tests to ensure the new intrinsic is called appropriately and that no `call_extern` is used for internal TileLang operations.
- Enhanced documentation in SKILL.md regarding the addition of new TileLang intrinsics.
* Add bfloat16 hexp overload and update common.h for compatibility
- Introduced `hexp` overload for `bfloat16_t` in `common.h` to bridge TileLang and CUDA's handling of exponential functions.
- Updated tests to verify the presence of the new `hexp` definition in `common.h`.
- Enhanced CUDA code generation to support the new intrinsic for better performance and compatibility.
* Update common.h for CUDA compatibility with aligned double4 types
- Adjusted preprocessor condition in `common.h` to maintain compatibility with CUDA versions prior to 13 by ensuring proper definition of aligned double4 types.
- This change enhances the compatibility of the code with older NVRTC built-in vector headers.
* Refactor kernel launch materialization
* [Refactor] Improve git commit ID retrieval and version labeling logic
* Enhance the get_git_commit_id function to check for the existence of the .git directory before attempting to retrieve the commit hash, ensuring more robust behavior.
* Update the dynamic_metadata function to conditionally set the NO_VERSION_LABEL environment variable based on the presence of the .git directory.
* Modify the GitHub Actions workflow to unset the NO_VERSION_LABEL environment variable during the installation and version check of the package, allowing for accurate version labeling in SDist builds.
* [Refactor] Update WebGPU backend registration
* Import the WebGPU package to automatically register its pass pipeline.
* Remove the explicit registration of the WebGPU pass pipeline, streamlining the backend setup process.
* Place auto WS producers in first warp group
* [Transform] Update layout inference to bind thread range for fully replicated fragments; add tests for TMA fragment mask copy and producer-consumer warp specialization
* [Refactor] Update thread variable calculation to use zero-based indexing across GEMM operations
* [Transform] Enhance simplification process with ContextSingletonFloorDivFolder and add tests for floor division simplification
* [Refactor] Simplify tensor conversion in dequantization functions and improve performance in GEMM examples
[BugFix][Transform] Freshen buffer defs after loop unroll
UnrollLoop clones each loop body with Substitute, which preserves Buffer object identity for definitions inside the body. This can define the same DeclBuffer handle in multiple unrolled copies after flattening and violate TIR SSA. Freshen buffer definitions per cloned body while preserving DeclBuffer backing data aliases.
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* [Backend] Refactor Pipeline to support different backends
Introduce Pipeline abstraction in backend/pipeline.py with per-backend
registration. Each backend (cuda, hip, c, llvm) now registers its own
compilation pass pipeline. engine/lower.py resolves the pipeline from
the target instead of hardcoding phase imports.
* enhance the structure
* fuse two stage
* fix after rebase
* update template pass pipeline
* pipeline
* fix test
* fix test
* fix cuda pipeline prologue
Ruinan Li