* [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>
Fix MFMA dtype args passed as DataType instead of str
When local_size is 1, compute_a_dtype/compute_b_dtype/compute_out_dtype
are assigned directly from self.a_dtype etc., which are tvm_ffi DataType
objects (a str subclass). The TVM FFI C++ layer dispatches DataType before
str in its type conversion, so these never get auto-converted to StringImm
(a PrimExpr). This causes MFMA GEMM compilation to fail on ROCm with:
TypeError: Expected Array<ir.PrimExpr> but got Array[index 3: DataType]
The bug only triggers when local_size == 1 (e.g. float32 GEMM), because
when local_size > 1 the f-string produces a plain str.
Fix: convert dtype attrs to str() before building compute dtype strings.
Add test: testing/python/amd/test_mfma_dtype_str_bug.py
* [feat] add lower-trace support for debugging
* [fix] address review: rebase, doc migration, lower_trace robustness
Address CodeRabbit + SiriusNEO review feedback on #2470:
Blocking (SiriusNEO):
- Rebase onto tile-ai/main (resolve docs/ + tilelang/__init__.py conflicts).
- Move IR Lower Trace doc from docs/tutorials/ to docs/tools/lower_trace.md
(alongside pass_diff.md); add index entry and 'superseded' pointer on
pass_diff.md; cross-link from the tutorial.
CodeRabbit Major:
- lower_trace/__init__.py: finally block uses active_exc pattern so a
generate_html() failure can no longer mask the real lowering exception.
- lower_trace/html.py: sanitize rec.phase into DOM-safe IDs via _safe_id()
before embedding in id="..." / selectors (avoids markup/JS breakage and
local HTML injection from special characters).
CodeRabbit Minor:
- test_lower_trace.py: restore caller's TL_LOWER_TRACE* env after the guarded
tilelang import instead of permanently clearing it.
- lower_trace/__init__.py: terminal mode prints a FAILED stanza (pass name,
exception, pre-pass IR) before re-raising so the failing step is visible.
- lower_trace/core.py: _get_codegen_output_path() no longer auto-writes
codegen.cpp in terminal-only mode (explicit override still honored).
- lower_trace/html.py: expandAll()/collapseCtx() null-guard the table lookup
so a changed diff with no table fails safe.
Cleanup:
- Add docstrings to test helpers and tools.__getattr__ (coverage now 94%+).
- ruff check + format clean.
* [Docs] Add lower trace to tools sidebar
* [fix] route lower-trace env vars through tilelang.env
Register TL_LOWER_TRACE / TL_LOWER_TRACE_DIR as EnvVar in the central
Environment class with get_lower_trace_mode()/get_lower_trace_dir()
parsers, mirroring the existing TILELANG_PASS_DIFF convention, instead
of reading os.environ directly in the lower_trace module.
* [fix] simplify lower_trace lazy-import in tools __getattr__
Drop the redundant globals() cache assignment (import_module already
binds the submodule onto the parent package) and fix the docstring: the
module is imported lazily for cost, not to defer monkey-patch side
effects, which happen in enable(). Note why importlib is required over
"from . import" (the latter re-enters __getattr__ and recurses).
---------
Co-authored-by: caojian5 <caojian5@huawei.com>
Co-authored-by: SiriusNEO <chaofan@deepseek.com>
* [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>
