common.h defines a bfloat16_t __hfma wrapper that forwards to the
__nv_bfloat16 __hfma overload. CUDA only declares that overload for
__CUDA_ARCH__ >= 800, so on SM75 merely parsing the wrapper fails with
"no instance of overloaded function __hfma". Because common.h is included
by every generated kernel, this breaks all SM75 builds, not only bf16 ones.
Guard the native path on SM80+ and use an fp32 FMA followed by bf16
conversion on earlier architectures, following the existing pre-SM80 bf16
fallback strategy.
Add a compile-only regression test that builds a bf16 ieee_fmaf kernel for
an explicit SM75 target. Existing IEEE math tests use target="cuda", so an
SM80+ runner exercises only the native branch.
* Add language dialect overlays
* Refactor language dialects into static backend layers
* Make language facade a static re-export and extract sparse layout leaf
Replace the runtime-activated language facade with a static
`from tilelang.cuda.language import *` re-export: drop the .pyi stubs +
generator + py.typed, the globals()-based __all__ scraping, and the
_activate_cuda_facade() setattr. Backend dialects (cuda/metal/rocm) now
build on tilelang.language.common with per-backend tir overlays.
Import-time dtype defaults in the mma/wgmma/mfma macro generators are
pinned to the dtypes leaf so they no longer touch the half-initialized
facade during bootstrap. Metal generators import tilelang.metal.language
(not the cuda facade) so the eager builder's real getattr resolves the
simdgroup ops; a lazy __getattr__ breaks the metal.language <-> intrinsics
cycle.
Extract the 2:4 sparsity layout metadata into
tilelang/cuda/intrinsics/sparse_layout.py (a dtypes-only leaf) so mma_sp
imports it directly and tilelang.utils.sparse drops out of the bootstrap
closure, letting its @tilelang.jit compressor go back to a module-level
decorator.
Exports unchanged: common=558, cuda/facade=683.
* Remove unwired BackendModule scaffolding
The BackendModule manifest and per-backend backend.py registrations were
included by mistake: nothing imports the per-backend backend.py files, the
ensure_*_loaded lazy-loaders have no call sites, and the mechanism
duplicates the register_lazy_execution_backends wiring already in
tilelang/backend/__init__.py. Drop the dead code and its unit test.
* Import tvm.arith in metal macro generator
The simdgroup stride check uses arith.Analyzer but the module was missing
the tvm.arith import.
* Use string dtype defaults in CUDA macro generators
Default dtype arguments are evaluated at import time, when the language
facade is only partially initialized, so they cannot reference T.*. They
previously used the tilelang.language.dtypes leaf; since every consumer
treats these as strings (annotated str, wrapped via DataType before use),
plain string literals are equivalent and drop the _dtypes dependency.
* Drop MPSIntrinEmitter from metal language facade
The emitter is only consumed directly from tilelang.metal.intrinsics, never
via the language surface, and exposing it here required a lazy __getattr__ to
break the metal.language <-> metal.intrinsics import cycle. Removing it from
__all__ eliminates the cycle and the deferred import.
* Restore dtype forwarding for ptx_ldmatrix
T.ptx_ldmatrix previously accepted a dtype= keyword via _dtype_forward (from
the old language/ast/ir.py path). The CUDA dialect alias dropped the wrapper
while every neighboring ptx_* operator kept it, changing the public call
contract. Wrap it to match.
* Re-export tir Select/min/max wrappers and dedup metal __all__
Add Select/min/max to tilelang.language.tir.ir __all__ so the common
facade re-exports the documented wrappers instead of the raw TVM
symbols. Build metal __all__ via dict.fromkeys to match the sibling
dialects and satisfy PLE0604.
---------
Co-authored-by: SiriusNEO <chaofan@deepseek.com>
* [BugFix] Transpose only the final two axes
* [Test] Cover transpose boundary predicate mapping
* [BugFix] Guard transpose bounds with generated indices
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* [BugFix] Reject contracting shared-buffer layouts in T.annotate_layout
T.annotate_layout could accept a contracting many-to-one forward map on a
shared buffer. makeBufferWithLayout then misread the smaller output extent as
a replication factor, rewrote the buffer to a higher rank, and eventually
failed with an unrelated BufferStore dimensionality ICHECK.
Add a pigeonhole check for static layout shapes before the replication branch.
When the layout output extent is smaller than its input extent, raise a
user-facing ValueError that identifies the buffer and reports both extents.
Keep identity, permutation, and output-growing layouts unchanged.
This check intentionally targets contracting maps; general injectivity
analysis for maps with larger output bounding boxes is outside this change.
Fixes#2632
* [Layout] Enforce injective shared-buffer annotations
Add a LayoutNode injectivity check with an exact fallback for static domains, and reuse it for FragmentNode thread and replication coordinates. Validate shared-buffer annotations after layout reshaping so sparse many-to-one maps are rejected with a useful collision diagnostic while padded injective maps remain valid.
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* [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.
