* [BugFix] Pack 32-lane 8-bit CUDA vectors correctly
* [Test] Avoid SM100-only store in int8 broadcast regression
* [BugFix] Move int8x32 packing into CUDA templates
* [BugFix] Track GEMM accumulator writes in warp-specialization liveness collector
* [BugFix] Track scale-factor regions and region extents in gemm liveness
T.atomic_load and T.atomic_store previously accepted every order in the shared memory-order map, including operation-illegal combinations that trigger device-side assertions in libcu++.
Validate memory orders against operation-specific allow lists before lowering, report invalid values with a frontend ValueError, and update the documentation and regression coverage accordingly.
Fixes#2574
* [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>
`T.infinity` aborted with an internal `LOG(FATAL)` ("Cannot decide
infinity for type float8_e5m2") because `infinity_op` had no fp8 branch.
e5m2 is the only fp8 format with a representable inf, so add a branch
returning +inf, and emit it in codegen via its bit pattern
(`::bitcast(0x7c)`, mirroring the tfloat32 path) since inf/NaN cannot be
spelled as a float literal for the fp8 constructor. Other fp8 formats
(e4m3*, e3m4, e8m0, finite fnuz incl. e5m2fnuz) have no inf and keep the
fatal. Adds e5m2 to the infinity test.
* feat(autotuner): support per-config pass_configs with code deduplication
- Extract _merge_pass_configs_into_compile_args helper to deduplicate
pass_configs merge logic across grouped_compile.py and tuner.py
- Replace json.dumps uniformity check with direct equality comparison
- Extract _make_jit_compile_func factory to unify lazy/eager jit_compile
- Fix decorator-mode cache isolation: use compile() directly when
per_config_pass_configs is present to avoid JIT cache key collision
- Add example and test for per-config pass_configs tuning
* docs: add docstrings to meet coverage threshold
* refactor(autotuner): move pass_configs resolution to tuner so grouped compile handles uniform configs
* fix(autotuner): define _PASS_CONFIGS_KEY constant and fix jit_compile race condition
- Replace all magic string "__pass_configs__" with _PASS_CONFIGS_KEY constant
- Move _pass_configs_lock to cover both branches in jit_compile to prevent
concurrent reads of temporarily mutated pass_configs
* refactor(autotuner): introduce type aliases for compile-unit bookkeeping
- Add ConfigArg, UnitItem, UnitResult, CompileUnit type aliases
- Replace all long tuple type annotations with aliases
- Expand example configs to include num_stages=[1,2] and enable_grouped_compile
---------
Co-authored-by: daniu.ld <daniu.ld@alibaba-inc.com>
T.dp4a previously accepted arbitrary buffer dtypes and passed them to backend DP4A implementations, which could silently reinterpret the operand bytes and produce incorrect results.
Require both inputs to use int8 and the accumulator to use int32. Add regression coverage for valid Buffer and BufferLoad forms and for invalid input and accumulator dtypes.
Fixes#2600
* Add experimental IKET CUDA backend hooks
* Keep IKET APIs under namespace
* Use real Perfetto screenshot for IKET example
* Show IKET events in Perfetto timeline screenshot
* Add IKET validation overhead example
* Remove IKET validation overhead example
* Fix IKET lint issues
* Move IKET integration to CUDA tools
* Add tools documentation and fix workflow regressions
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* [Feature] Add pass_visualizer: structure-tree pass browser
Add an interactive, pass-by-pass IR structure-tree visualizer under
tilelang/utils/pass_visualizer. It complements the existing text-level
pass_diff tool by rendering the SBlock structure tree and expanding tile
ops by field name, with per-class operator highlighting (tile op / sync /
lowered hardware intrinsic) across the CUDA lowering prologue.
Includes a CLI entry point, an example gemm_relu kernel, a pytest suite,
and a docs section contrasting it with pass_diff.
Co-Authored-By: Claude <noreply@anthropic.com>
* Move pass_visualizer to tilelang/tools and fix prologue + review issues
- Move package from tilelang/utils to tilelang/tools per review
- Add missing MaterializeKernelLaunch pass (fixes CUDA CI modulo-by-zero)
- Escape "</" in embedded JSON to prevent <script> injection in HTML report
- Use explicit utf-8 encoding for source/output file I/O
- Guard _fmt_shape against symbolic shapes and validate module spec
Co-Authored-By: Claude <noreply@anthropic.com>
---------
Co-authored-by: Claude <noreply@anthropic.com>
* [CUDA] Add T.__fns intrinsic for find-nth-set bit
Expose CUDA __fns in TileLang for direct lookup of the k-th active lane in a 32-bit bitmask, complementing the existing T.__ffs helper.
* Add builtin intrinsics for shared ldst and atomics
* [BugFix] Skip source-compilation options when exporting LLVM module cache
* Add regression test
* Refactor kernel cache export handling and improve test coverage
- Renamed `_get_compile_args` to `_get_source_compile_args` for clarity.
- Introduced `_get_export_link_args` to manage export link arguments.
- Updated `_safe_write_executable` to accept export kwargs directly.
- Enhanced `TVMFFIKernelCache` to determine export kwargs based on target host.
- Modified tests to validate new export behavior and ensure source options are correctly handled for different targets.
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
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.
