Use PTX v4 atomics for 16-bit CUDA atomic_addx4
The initial fp16/bf16 atomic_addx4 fix avoided the generic float4 path by splitting each operation into two AtomicAddx2 calls. That fixed correctness, but it still used the simulated implementation on Hopper where PTX provides native vector atomics.
Add sm90+ half_t and bfloat16_t AtomicAddx4 overloads backed by the PTX vector forms atom.global.v4.f16.add.noftz and atom.global.v4.bf16.add.noftz, including the existing memory-order variants. Keep the x2 fallback for pre-sm90 targets because PTX vector f16/bf16 atomics require sm90 or newer.
Keep the CUDA language regression focused on fp16/bf16 atomic_addx4 codegen and offset coverage.
Co-authored-by: dingsg <shengge.ding@enflame-tech.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.
_tir_u32_to_int_to_float decoded signed sub-word fields from uint32 storage as unsigned values because it only masked the selected field before casting. As a result, negative encodings such as signed int4 `0xF` decoded as `15.0` instead of `-1.0`.
Sign-extend the extracted field through int32 before casting to the requested float dtype, matching the existing packed signed-int decode helper.
Add CUDA regression coverage for signed int2, int4, and int8 values decoded from uint32 storage.
Fixes#2482
* Add split-k for example_warp_specialize_flashmla
* Rotate QK issue to end of main loop to allow more latency hiding with TMA issue
* During rescale, read max score from REG instead of SMEM
* Move K_pe_shared_1 load
* Use stmatrix for SP0 write to reduce SMEM bank conflict
* Support arbitrary SMEM layout in MMA lowering
* Remove runtime guard for increase_descriptor_offset
* [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.
* [BugFix] Reject unsupported T.serial access to thread-distributed fragments with clear diagnostics
A T.serial loop runs sequentially inside each CUDA thread, but a fragment's
elements are distributed across threads, so indexing a fragment by a serial
loop variable has no valid thread-ownership mapping. Three fuzzer reports share
this root cause:
- #2393: serial reduce of a GEMM-output fragment -> unhelpful internal error
("contains inner var j").
- #2395: serial write into a fragment then T.copy out -> nvcc "identifier
undefined" (an owner-thread guard built from serial loop vars was emitted
outside the loop that binds them).
- #2396: serial write of a fragment to global -> silently wrong (threads
race-write each cell).
AddWrapperForSingleBufStore wraps bare fragment stores in a degenerate
T.parallel(1) and only rejected constant non-zero fragment indices, so a
dynamic (e.g. serial-loop) index slipped through and lowered to broken code
(#2395/#2396). Reject any fragment index that is not the constant all-zero
index in this fallback path, with a message pointing to T.Parallel /
T.reduce_sum. Rephrase the existing inner-var layout check (#2393) to explain
the cause instead of dumping "contains inner var".
Also collect every access to a buffer in a statement (not just the last) so the
diagnostic is not bypassed when one fragment is accessed at both a constant and
a variable index in the same statement; without it such a case falls through to
a cryptic downstream StructuralEqual internal error. This is a completeness
hardening only, not required by the three issues above.
Fixes#2393Fixes#2395Fixes#2396
* Generalize fragment diagnostic to non-parallel loops
inner_vars_ is populated for every non-parallel For (serial, unroll,
vectorized), so the message must not hard-code T.serial. Reword to
"non-parallel loop variable" per review feedback.
* Narrow fix scope to #2393 only
Remove fragment index validation from AddWrapperForSingleBufStore and
delete tests for #2395/#2396. These issues cannot be caught at this pass
stage because T.grid and T.serial are indistinguishable (both appear as
ForKind.SERIAL before LowerTileOp).
Keep only the ParallelOpNode diagnostic improvement which successfully
catches #2393.
* Add CUDA target code attr support
* Add CUDA target normalization support
Introduced a new function to normalize CUDA targets, ensuring that the target is correctly identified and transformed based on the detected architecture. Added a test to verify that the bare CUDA target uses the detected architecture accurately.
* Add cross-host CUDA binary cache
* Nest host kernel cache under version namespace
* [BugFix] Fix T.Persistent dropping tiles when last dim is not a multiple of group_size (#2433)
* Update src/ir.cc
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
---------
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
* Add CUDA target code attr support
* Add CUDA target normalization support
Introduced a new function to normalize CUDA targets, ensuring that the target is correctly identified and transformed based on the detected architecture. Added a test to verify that the bare CUDA target uses the detected architecture accurately.
T.fill of an int8 buffer with a negative scalar (other than -1) silently
wrote only 1 of every 4 elements correctly and clobbered the other 3 to
-1. The vectorized fill packs four int8 lanes into one 32-bit store via
make_int(v,v,v,v); a negative signed char is integer-promoted to int and
sign-extended before the shifts/OR, so its high 1-bits flood the
neighbouring lanes (e.g. make_int(-7,-7,-7,-7) returned 0xFFFFFFF9 =
[-7,-1,-1,-1]).
Build the packed value from explicit unsigned bytes so the sign-extension
bits cannot leak across lanes. Add an int8 negative-fill regression test.
Fixes#2427
* Add stmatrix m16n8 support
* lint
* Refactor stmatrix code generation to handle optional shape argument
Updated the handling of the shape argument in the stmatrix function across CUDA, CuTeDSL, and HIP code generators. The shape is now checked as the last argument, and the payload start index is adjusted accordingly. This change improves consistency and clarity in the code generation process for tile language operations.
* Refactor stmatrix function naming for consistency
Updated the naming convention of stmatrix functions in CUDA and related files to improve clarity and consistency. The function names now follow a uniform pattern, placing the shape descriptor before the axis identifier. This change enhances readability and aligns with the recent updates in code generation logic.
* Enhance extern call statement printing in CUDA codegen
Updated the `VisitExpr_` method in `codegen_cuda.cc` to improve the clarity of extern call statements. The function now includes comments indicating that the transformation, number, and shape are encoded in the helper name, and only the shared memory pointer and value operands are printed. This change aims to enhance code readability and maintainability.
* Refactor extern call expression printing in CUDA codegen
Introduced a helper function to streamline the printing of extern call expressions in the `VisitExpr_` method of `codegen_cuda.cc`. This change enhances code readability by reducing redundancy and improving the clarity of the generated output for specific function calls, such as `tl::pack_b8x4`.
* Refactor shape argument handling in stmatrix code generation
Updated the handling of the shape argument in the `VisitExpr_` method for both CUDA and CuTeDSL code generators. The shape is now checked as the fourth argument instead of the last, improving clarity and consistency in extern call statement printing. This change enhances the robustness of the code generation process for tile language operations.
* Refactor shape argument handling in stmatrix code generation
Updated the handling of the shape argument in the `VisitExpr_` method for both CUDA and CuTeDSL code generators. The shape is now checked as the last argument and validated accordingly, improving clarity and consistency in the code generation process for tile language operations.
* Refactor tilelang lower function context management
Updated the context management in the `test_tilelang_copy_uses_stmatrix_m16n8_for_sm100_int8_shared_store` test by removing the unnecessary `PassContext` wrapper. This change simplifies the code and enhances clarity in the test setup for tile language operations.
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
Vectorized atomic_add into an fp16/bf16 destination with an fp32 source
bit-reinterpreted the source float2 as half2/__nv_bfloat162 instead of
numerically converting it, silently producing wrong results. This hit
both the auto-vectorized T.atomic_add path (AtomicAddx2 value form) and
the explicit T.atomic_addx2 API (pointer form).
Add float2 / float* converting overloads (round-to-nearest) to ToHalf2
and ToBfloat162 so the vectorized helper converts instead of
reinterpreting the source bytes. Add mixed-dtype regression tests for
both entry points.
Fixes#2382
feat: add end-to-end LLVM backend execution support
Enable LLVM as a TileLang execution backend in CMake and route llvm targets through TVM LLVM codegen during lowering.
Fix tvm_ffi output tensor device inference by deriving the output device from tensor inputs, and add LLVM execution regression tests for matmul, T.gemm, T.copy, and T.While.
A vectorized T.copy between fp16 and bf16 failed to compile. The CUDA
cast codegen has dedicated vectorized branches for fp16<->fp32,
bf16<->fp32, fp8 and fp4, but none for fp16<->bf16, so it fell into the
elementwise fallback. There the lanes load as native __half/__nv_bfloat16,
and a direct (half_t)(__nv_bfloat16) (or reverse) is an ambiguous
user-defined conversion, breaking the build.
Route cross-half elementwise casts through float to disambiguate, and
add a cross-dtype copy regression test for both directions.
Fixes#2385
* [BugFix] Complete CuTeDSL scan lowering for TileKernels
TileKernels exercises CuTeDSL scan and routing patterns that were not fully covered by the existing Python backend helpers.
The MoE get_fused_mapping kernel computes expert base offsets with T.cumsum over a 256-entry shared buffer. CuTeDSL CumSum1D only performed a single warp-local prefix scan, while the CUDA scan template uses the first warp to scan all 32-element segments and carry the result between segments. As a result, expert prefixes reset at later warp-sized chunks, positions for experts such as 32 and 64 overlapped earlier expert ranges, and pos_to_expert validation failed on H100.
Bring the CuTeDSL scan helpers in line with the CUDA InclusiveScanLine behavior: add reusable one-warp line scans with carry across 32-element segments, route CumSum1D and CumMax1D through them, and use the same line scan for CumSum2D/CumMax2D on both row-wise and column-wise axes. This removes the previous dim=0 H<=32 limitation and fixes multi-segment cummax as well as cumsum.
The same TileKernels path also relies on CuTeDSL codegen accepting non-flat buffer accesses after TIR lowering. Add a shared index linearization helper for loads, stores, address_of/reinterpret, and atomic pointer generation, and lower tl.match_any_sync through the existing CuTeDSL warp intrinsic. Thread binding now falls back to the IterVar name when thread_tag has already been normalized away.
Verified on H100 CUDA/CuTeDSL with CUDA_VISIBLE_DEVICES=0 and TILELANG_TARGET=cutedsl:
- cmake --build build -j$(nproc)
- ruff format --check tilelang/contrib/cutedsl/reduce.py
- ruff check tilelang/contrib/cutedsl/reduce.py
- git diff --check
- pytest testing/python/target/test_tilelang_codegen_cutedsl_scan.py testing/python/language/test_tilelang_language_scan.py -q -s --tb=short: 13 passed
- pytest testing/python/language/test_tilelang_language_warp_vote.py::test_match_any_sync -q -s --tb=short: 1 passed
- pytest tests/moe/test_get_fused_mapping.py::test_get_fused_mapping[num_send_tokens=4001-num_topk=2-num_experts=72-num_ep_ranks=1-alignment=64] -q -s --tb=short: 1 passed
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Resolve CuTeDSL dynamic shapes from live tensors
CuTeDSL records dynamic symbolic dimensions in the same ordering as the CUDA wrapper, but it previously kept only the first buffer shape or stride location for each symbol. That is not sufficient for kernels with optional tensor arguments. TileKernels reduce_fused declares num_tokens on topk_weights, token_topk_to_pos, and out; when with_weights is false the topk_weights argument is None, so the first recorded source is not a live tensor.
The previous adapter fallback substituted 0 when that first source was None. On H100 with TILELANG_TARGET=cutedsl this passed num_tokens=0 to the generated launcher even though token_topk_to_pos had shape (4001, 2), producing grid=[0, 1, 1] and a CUDA error 1 launch failure in reduce_fused.
Keep all shape/stride candidates for each dynamic symbol and resolve the runtime value from the first candidate backed by a real torch.Tensor. Reuse the same resolution helper for allocated output shapes and for dynamic arguments passed to the generated CuTeDSL module, and fail clearly if no live tensor source exists.
Verified with a CuTeDSL adapter regression test for repeated dynamic shape symbols with a None first candidate, and with the TileKernels reduce_fused H100 failure case that previously launched with num_tokens=0.
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Permit dead optional CuTeDSL stride symbols
TileKernels no-scale MoE expand paths can keep a dynamic stride symbol in the generated CuTeDSL host wrapper ABI even when the optional scale tensor is absent from the lowered device kernel variant. In that case the runtime argument is intentionally None, so the stride-only symbol has no live tensor from which the adapter can read torch stride metadata.
The previous dynamic-symbol fix made missing live tensor sources strict to avoid unsafe shape fallbacks such as launching with num_tokens=0 when the first candidate tensor was optional. That strict behavior is still required for shape symbols because they control output allocation and launch dimensions.
Limit the fallback to symbols whose candidates are stride-only. Such values are ABI placeholders for optional strided tensor parameters in pruned variants; returning 0 preserves the generated call signature without weakening dynamic shape resolution. Add adapter regressions covering both the optional stride fallback and the strict missing-shape error.
Verified on H100 with TILELANG_TARGET=cutedsl:
- ruff format/check on the touched files
- pytest testing/python/jit/test_tilelang_jit_cutedsl_host_codegen.py::test_cutedsl_adapter_resolves_dynamic_symbol_from_live_tensor_candidate testing/python/jit/test_tilelang_jit_cutedsl_host_codegen.py::test_cutedsl_adapter_allows_optional_stride_symbol_without_live_tensor -q
- pytest tests/moe/test_expand_to_fused.py::test_expand_to_fused[num_send_tokens=4001-num_topk=2-num_experts=9-num_ep_ranks=8-hidden=576] -vv -s --tb=short
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Fix CuTeDSL optional ABI and early-return codegen
TileKernels full functional coverage with TILELANG_TARGET=cutedsl exposed additional CuTeDSL gaps in top2_sum_gate after the broader TileKernels fixes.
First, the host adapter still appended dynamic shape symbols for optional tensors that had been compiled out by static flags. In top2_sum_gate, to_physical_map/logical_count can be None when logical expert mapping is disabled, but their shape-only symbol remains in the wrapper ABI. Keep output allocation strict, but allow dead host ABI dynamic shape arguments to use a zero placeholder when no live tensor source exists.
Second, CuTeDSL cannot emit dynamic thread returns directly, so codegen rewrites if cond: thread_return(); rest into if not cond: rest. The existing recognizer only handled a bare thread_return then body, and missed the common if cond: stores; thread_return() pattern. Preserve the pre-return stores and guard the following statements so masked tokens cannot be overwritten by the fallthrough path.
Third, vector binary codegen reused SSA aliases for operands containing BufferLoad nodes. That is not valid for mutable local tensors after an intervening store. The softmax top2_sum_gate path restored raw logits to scores_local, but the subsequent scores_local + bias expression reused an older softmax load and ranked by softmax(logits)+bias instead of logits+bias. Avoid SSA alias reuse for vector operands containing BufferLoad so the generated code reloads the current tensor value.
Verified on H100 with the local TileLang build and TileKernels CuTeDSL backend:
- cmake --build build -j 32
- ruff format/check on touched Python files
- pytest testing/python/jit/test_tilelang_jit_cutedsl_host_codegen.py::{dynamic symbol candidate tests} -q: 3 passed
- tests/moe/test_top2_sum_gate.py::test_top2_sum_gate[num_groups=0-num_topk_groups=0-num_routed_experts=72-num_shared_experts=1-num_topk=6] with TILELANG_TARGET=cutedsl TK_FULL_TEST=1: 1 passed
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Promote narrow CuTeDSL shifts before wide casts
TileKernels e5m6 quantization kernels pack 8 truncated fp16 values into three uint32 words. The TileLang source writes expressions such as T.cast(half_u16[i] << 20, T.uint32), which the CUDA C++ backend evaluates with normal integer promotion before the final uint32 cast.
CuTeDSL uses strongly typed cutlass integer wrappers instead of CUDA C integer promotion. The previous CuTeDSL lowering cast the shift result back to the narrow TIR result type first, producing code like cutlass.Uint32(cutlass.Uint16((half_u16[0] << cutlass.Uint16(20)))). For shifts by more than the uint16 width this drops all high bits before the packed uint32 store, corrupting e5m6 payload bytes. Cast-back then faithfully decoded the corrupted payload into nan/inf values, so TileKernels full correctness failed in cast_back_e5m6 and per_token_cast_to_e5m6 cases.
When a narrow integer shift is immediately cast to a wider integer type, emit the CuTeDSL shift with the left operand promoted to the target integer type before shifting. This preserves the intended packed-bit construction while leaving narrow shift expressions that are not widened on the existing path.
Add a CuTeDSL codegen regression test for uint16 << 20 widened to uint32 so future changes do not reintroduce a narrow in-place shift.
Verified on H100 CUDA/CuTeDSL with TILELANG_TARGET=cutedsl and TILELANG_DISABLE_CACHE=1:
- cmake --build build -j 32
- python -m pytest testing/python/target/test_tilelang_codegen_cutedsl_integer_shift.py -q
- TileKernels focused e5m6 quant tests: tests/quant/test_per_token_cast_to_e5m6.py tests/quant/test_cast_back_e5m6.py -n 2: 168 passed, 112 skipped
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
---------
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
`hexp` is #define'd to `cutlass::fast_exp`, so the hand-written
`cutlass::fast_exp(bfloat16_t)` overload's `return ::hexp(x)` expanded into a
call to itself: bf16 `T.exp` recursed and returned garbage.
Route through float, mirroring cutlass's own `half_t fast_exp`
(cutlass/fast_math.h), so the call resolves to `cutlass::fast_exp(float)`
(= ::expf). Verified against torch.exp using the #2383 repro.
Fixes#2383.
* Fix TMA 1D load test lowering
* Refactor test execution in tilelang TMA 1D test script
Updated the main execution flow to directly call the test function instead of running an elementwise add operation, improving clarity and focus on testing functionality.
The SM75 GEMM emitter derives MMA shapes from input dtype bits alone, so it
already emits the (16,8,8) shape for fp16->fp16 accumulation (identical to the
registered fp16->fp32 case) and the (8,8,16) shape for uint8 inputs (identical
to signed int8). But the C++ dispatcher table only covered the three configs
tested by #2198 (fp16->fp32, int8, int4), so requesting fp16->fp16-accum or
uint8 GEMM on Turing failed at nvcc with an opaque
`tl::mma_sync: unsupported configuration` static_assert.
Add the two missing tl::detail SM75 atoms and register their dispatchers,
following the precedent set by #2198's hand-written int4/uint4 structs (CuTe
ships only fp16->fp32 and signed-int8 SM75 atoms). The uint8 atom uses the
non-saturating PTX form to match the existing signed-int8 path.
Verified on a TITAN RTX (SM75): emitted PTX is
mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 and
mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32, and end-to-end GEMM numerics
match torch.
* fix: name
* [Transform] Enhance role management and scalar dependency propagation
- Introduced a new Role enumeration value `kNeutral` to better represent variable states.
- Added `CombineRoles` function to manage role combinations effectively.
- Implemented `ScalarAccess` structure and `ScalarAccessCollector` class for tracking variable definitions and usages.
- Enhanced `PropagateScalarDependencyRoles` function to propagate roles based on scalar dependencies across statements.
- Updated role assignments in various statement visitors to utilize the new role management logic, improving clarity and correctness in role handling.
* Update main execution flow in test_tilelang_language_let.py to enable tilelang testing framework
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* 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.
* [BugFix] CuTeDSL backend support for TileKernels benchmarks
TileKernels benchmark coverage with TILELANG_TARGET=cutedsl exercises backend paths that were either missing or not legal for CuTeDSL-generated Python/MLIR code. The failures were not limited to one kernel: packed FP4/FP8 quantization flows need sub-byte vector materialization, MoE mapping uses warp vote/shuffle and grid synchronization, and benchmark launch wrappers need dtype-aware host plumbing for generated CuTeDSL modules.
Before this change the H100 benchmark run exposed several classes of failures: missing shuffle/vote intrinsics such as shfl_xor/match_any/popc in CuTeDSL lowering, invalid scalar/vector casts for FP4 E2M1 and packed UE8M0 paths, dynamic if expressions whose branches narrowed or changed integer types, unsupported sub-byte load/store/copy patterns, and a fragile one-counter software grid barrier that could reset before all waiting blocks observed release.
Lower the missing TileLang operations to CuTeDSL runtime helpers, add Float4E2M1FN software pack/unpack and zero-fill support, make sub-byte vector loads/stores use packed Uint8 storage where required, cast conditional branches to the expected result dtype, and teach the CuTeDSL wrapper about float4/float8 dtype plumbing and C++ min/max expression rendering. Also replace the sync_grid single-counter barrier with a counter plus generation sense barrier so a waiting block cannot miss the release when the arrival counter is reset.
This keeps the fixes in the CuTeDSL backend rather than adding TileKernels-specific workarounds. For example, TileKernels expand_to_fused_with_sf relies on TileLang safe-memory legalization for aligned SF loops; the CuTeDSL generated code now preserves the hidden_sf guard instead of requiring the benchmark source to special-case packed UE8M0 layouts.
Validation on H100 (sm_90, 114 SMs) with CUDA visible device 0, Python 3.12, TILELANG_TARGET=cutedsl, TK_FULL_TEST=1, TILELANG_DISABLE_CACHE=1, and PYTHONPATH=/home/jayce.su/new/tilelang:
- cmake --build build -j 16
- ruff check tilelang/contrib/cutedsl tilelang/jit/adapter/cutedsl tilelang/jit/adapter/utils.py
- git diff --check
- Original failing node: tests/moe/test_expand_to_fused.py::test_expand_to_fused_with_sf_benchmark[num_send_tokens=4001-num_topk=2-num_experts=4-num_ep_ranks=64-hidden=576-num_per_channels=32-use_tma_aligned_col_major_sf=True-round_sf=True-use_packed_ue8m0=True], 1 passed with -n 2
- Full-run local pressure slice around the failing MoE benchmark region, 80 passed with -n 2 in 379.53s
Author-email: jayce.su@enflame-tech.com
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Complete CuTeDSL support for TileKernels benchmarks
TileKernels benchmark coverage with TILELANG_TARGET=cutedsl exposes several general CuTeDSL backend gaps. These are not TileKernels-specific workarounds, H100-only specializations, or changes required only because of a particular CUDA/CUTLASS DSL version pairing. The H100 CUDA 13.2 + nvidia-cutlass-dsl-libs-base 4.5.0 environment is the reproducer that exercises the missing backend coverage; the underlying issues are generic to the emitted CuTeDSL code shapes.
Root causes:
- Predicated T.copy and safe memory legalization lower BufferLoad-or-zero values into conditional expressions. Emitting these values through tl.where evaluates both branches eagerly, which can issue unsafe packed loads and can force CuTeDSL/MLIR through unsupported vector select and unrealized_conversion_cast paths for FP8/FP4 storage.
- Safe memory legalization can bind the zero branch through a local zero-like variable, let binding, cast, or broadcast-zero tensor. The CuTeDSL codegen must recognize those zero-like values when deciding whether a conditional load can be lowered lazily and when materializing Float4E2M1FN broadcasts.
- CuTeDSL generated code has to distinguish rmem Tensor objects from TensorSSA values. Appending .load() blindly is invalid once a branch or cast already produced TensorSSA, and this showed up in the TileLang CI cast example as AttributeError: 'TensorSSA' object has no attribute 'load'.
- Scalar FP8 widening and narrowing require CuTeDSL tensor casts on aligned vector values. Plain scalar construction falls back to unsupported scalar FP8 conversion paths, so the backend has to bridge scalar casts through a packed vector form.
- Sub-byte vectors such as Float4E2M1FN cannot be copied lane by lane because CuTeDSL does not support sub-byte scalar dereference. The backend must preserve TensorSSA/packed-vector store form for these values, and scalarized local/shared sub-byte vector buffers must use the sub-byte packing factor as the make_tensor_at_offset div_by assumption.
- Generated CuTeDSL kernels can call CUDA-template math helpers such as __habs and __float2half_rz, and math wrappers such as log1p/isfinite. Those helper names must be provided by tilelang.contrib.cutedsl.math when the Python CuTeDSL backend is selected.
- sync_grid publishes block arrival through thread 0, but the memory ordering has to cover all threads in the block. A CTA sync alone does not provide device-scope ordering for global writes performed by non-zero threads.
Fixes:
- Recognize zero-like vars in conditional store lowering, including let bindings and bind values that are casted or broadcast-zero tensors. This keeps predicated loads lazy and avoids eager tl.where evaluation for packed and vectorized T.copy paths.
- Add CuTeDSL helper bridges for Tensor/TensorSSA normalization. tl.as_tensor_ssa avoids calling .load() on TensorSSA values, and tl.as_rmem_tensor spills TensorSSA values into rmem tensors before byte recasts for sub-byte stores.
- Add scalar FP8 bridge lowering that packs Float16/Float32/BFloat16 or FP8 scalar values into a four-lane rmem tensor, uses tl.cast_tensor, and reads lane zero. This matches CuTeDSL/NVGPU alignment requirements without forcing front-end kernels to perturb tile shapes just to avoid scalar casts.
- Keep sub-byte vector stores in TensorSSA form instead of copying lane by lane, keep non-load vector store RHS materialized as rmem tensors so CuTeDSL .store() receives a TensorSSA value, and use 8 / element_bits as the div_by packing factor for scalarized local/shared sub-byte vector buffer views.
- Gate the Float16 nvgpu.cvt_fpext fast path to floating-point source TensorSSA types only, so integer TensorSSA casts do not enter an FP-only conversion op.
- Strengthen sync_grid by fencing all block threads at device scope before thread 0 publishes arrival, then waiting at block scope and fencing again before user code resumes.
- Expose the missing CuTeDSL math helpers. __habs clears the sign bit for Float16/BFloat16 scalars and returns the original dtype. __float2half_rz uses inline PTX cvt.rz.f16.f32 so generated E5M6 cast-back code can preserve CUDA round-toward-zero semantics. log1p and isfinite are wired through the CuTeDSL math module and fastmath canonicalization.
- Keep the defensive dynamic-symbolic None fallback visible by logging a warning before substituting 0.
Version/specialization assessment:
- The patch is backend-generic: it fixes how TileLang emits CuTeDSL for predicated loads, FP8 scalar casts, Tensor/TensorSSA boundaries, sub-byte vector stores, grid synchronization, and generated math helpers. It does not specialize TileKernels cases, tensor shapes, H100 SM counts, or benchmark parameters.
- A future CuTeDSL version might accept some currently rejected IR forms, but relying on eager branch evaluation, missing helper symbols, incorrect TensorSSA .load() calls, or thread-0-only grid-sync ordering would still be fragile. The corrected lowering is semantically closer to TIR: predicated loads are lazy, sub-byte storage stays packed, TensorSSA values are not treated as tensors, and scalar FP8 conversion follows the vector cast path required by CuTeDSL/NVGPU.
- CUDA 13.2 and CUTLASS DSL libs 4.5.0 are recorded as the verified reproducer, not as the reason for the patch. The same TileLang backend gaps can be triggered by any workload that emits these CuTeDSL patterns.
Verified on H100 PCIe (sm_90, CUDA_VISIBLE_DEVICES=0 maps to NVIDIA H100 PCIe), CUDA 13.2, Python 3.12.3, PyTorch 2.12.0+cu132, nvidia-cutlass-dsl-libs-base 4.5.0, and local TileLang built from this branch:
- cmake --build build -j$(nproc)
- python3.12 -m py_compile tilelang/contrib/cutedsl/math.py tilelang/contrib/cutedsl/utils.py tilelang/contrib/cutedsl/grid_sync.py tilelang/jit/adapter/cutedsl/adapter.py
- pre-commit run --files $(git diff --name-only origin/main...HEAD)
- examples/cast/test_example_cast.py::test_example_group_per_split_token_cast_to_fp8: 1 passed
- representative per_token_cast, cast_back, cast_back_e5m6, per_block_cast_lossless, per_block_cast, and per_channel_cast benchmark failures: 6 passed
- TileKernels full benchmark with TILELANG_TARGET=cutedsl, TK_BENCHMARK_BACKEND=event, and pytest -n 2: 10165 passed, 160 warnings in 2:44:58
Jayce: JayceSu98 <jayce.su@enflame-tech.com>
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Avoid zero-sized fake tensors in CuTeDSL cubin generation
CuTeDSL cubin generation builds fake compact tensors to compile wrapper code before the real runtime tensors are launched. TileKernels exercises benchmark and correctness cases with empty dynamic dimensions, such as zero-token inputs. Passing those zero extents directly into make_fake_compact_tensor can make the CuTeDSL fake layout invalid even though the shape is only a compilation placeholder and the real runtime tensor metadata is passed separately.
Clamp only the fake tensor shape to positive extents before constructing the CuTeDSL placeholder. This keeps the cubin generation path valid for empty-input cases without changing runtime tensor shapes, strides, launch parameters, or kernel behavior.
This is not a TileKernels-specific workaround or a CUDA toolkit version workaround. It fixes the CuTeDSL adapter boundary so compile-time fake tensors satisfy CuTeDSL layout construction requirements while preserving the actual runtime contract.
Verified with:
- python -m py_compile tilelang/jit/adapter/cutedsl/wrapper.py
- git diff --check -- tilelang/jit/adapter/cutedsl/wrapper.py
Jayce: JayceSu98 <jayce.su@enflame-tech.com>
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
* [BugFix] Stabilize batched BF16 reduce test input
CUDA CI exposed a flaky failure in python/language/test_tilelang_language_reduce.py::test_reduce[sum-bfloat16-64x128-f2f-t256-b4]. The failing run had only one mismatched output element, with an absolute difference slightly above the 0.1 BF16 tolerance and a large relative error because the reference value was small.
The runtime code on this branch only changes CuTeDSL backend paths, and the same test passed locally on the ordinary CUDA backend when run in isolation. The failure is caused by the test input depending on the worker-local torch RNG state: the module seeds once at import time, but xdist load scheduling lets each worker consume a different number of random tensors before this parameterized case runs. That makes the BF16 sum case occasionally sample an input whose rounded reduction error crosses the tolerance boundary.
Make the parameterized reduce input schedule-independent by deriving a deterministic seed from the case parameters and using a local CPU torch.Generator for both integer and floating inputs. This preserves randomized coverage per case while decoupling correctness from xdist worker assignment, prior cases in the same worker, and full-suite ordering.
Verified on H100 PCIe with CUDA visible device 1, Python 3.12.3, PyTorch 2.12.0+cu132, and the local TileLang dev build:
- ruff check testing/python/language/test_tilelang_language_reduce.py
- python -m pytest -q -s --tb=short testing/python/language/test_tilelang_language_reduce.py::test_reduce[sum-bfloat16-64x128-f2f-t256-b4]
- python -m pytest --verbose --color=yes --durations=0 --showlocals --cache-clear --maxfail=3 --numprocesses=8 testing/python/language/test_tilelang_language_reduce.py
- 10 consecutive reruns of the failed BF16 reduce case
Jayce: JayceSu98 <jayce.su@enflame-tech.com>
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
---------
Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
LowerSharedBarrier lowers T.alloc_barrier() to hardware mbarrier code (the
cutlass Barrier type, tl::tl_shuffle_elect, tl::fence_barrier_init), all of
which only exist on sm_90+. On pre-Hopper targets (e.g. sm_86) the frontend
accepted the call silently and nvcc later failed with cryptic
"identifier 'Barrier' is undefined" / "namespace 'tl' has no member ..."
errors.
Add a static compute-capability guard in CUDAPassPipelineBody, before
LowerSharedBarrier: if the target is < sm_90 and the module uses a shared
barrier, raise a clear ValueError naming the offending primitives and pointing
to __syncthreads()/named barriers as the pre-Hopper alternative.
Fixes the sm_86 alloc_barrier compile crash.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* Fix eager JIT packed FP4 shape binding
* Add logical shape and stride calculations for buffers in eager builder
This update introduces helper functions to compute logical shape and stride values for buffers, enhancing the handling of packed storage factors. The changes streamline the process of determining dimensions and strides in the context of tensor operations, improving overall efficiency and clarity in the eager execution flow.
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* Add TMA copy for packed global FP4 to unpacked SMEM.
Wire IsValidTMACopyDtypePair and ALIGN16B tensor map descriptors for
float4_e2m1fn global to float4_e2m1_unpacked shared transfers, and
extend language tests for TMA and SIMT copy restrictions.
* Extend tcgen05 MMA metadata for FP4 f8f6f4 and dtype helpers.
* Refactor tcgen05_gemm_blockscaled API and lowering.
* Use separate A/B dtypes in GEMM tileop emitters across backends.
* Add maint/gemm correctness evaluation scripts and update gemm_sm100 examples.
* Fix profiler assert_allclose to compare outputs only.
* lint
* canonicalize data types
* add dpgm fp8-fp4 kernel
* resolve comments
* Refine FP4 unpacked SMEM dtype handling
Treat float4_e2m1_unpacked as a SMEM/TMA-only storage tag instead of a general vector/register dtype. Remove vectorized Python aliases and keep codegen uses constrained to shared-memory storage/reference paths that lower to uint8_t.
Allow packed FP4 global tensors to TMA-load into unpacked shared memory with the ALIGN16B TensorMap format, while keeping unpacked storage out of TMA store paths. Use payload-bit accounting for global strides and mbarrier transactions.
Add focused dtype and TMA descriptor tests, and validate the real tcgen05 fp4-unpacked GEMM paths.
* Restrict FP4 unpack TMA to loads
* Hide vectorized unpacked FP4 language aliases
* Update TVM for unpacked FP4 symbol exports
* Rebind TVM submodule to merged unpacked FP4 export fix
* Enable cache for persistent tcgen5 GEMM example
* 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
* [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>
