16-bit lowering emits half-style names for sinh, cosh, tanh, atan, erf,
and nearbyint, but no CUDA overload of those names accepts CUTLASS
half_t or bfloat16_t. Most toolkits do not declare the names at all, and
the overloads that do exist take native __half or __nv_bfloat16, which
the wrappers have no implicit conversion to, so these operations fail
everywhere.
Add hsinh, hcosh, htanh, hatan, herf, and hnearbyint overloads for both
wrapper types. Each evaluates in float32 and converts back, following
the existing htan bridge. This completes 16-bit bridge coverage for the
unary floating-point operations that CUDA lowering already handles.
Extend the 16-bit intrinsics test kernel to cover all six operations.
Introduce two CUDA-specific math intrinsics with guaranteed instruction
selection:
- T.fma(x, y, z): fused multiply-add whose fusion (single rounding) is
guaranteed regardless of NVCC contraction settings.
- T.fmul(x, y): multiply with an explicit boundary that is never
contracted into a neighbouring FMA.
Unlike the ieee_* family (which carries an explicit rounding mode and
stays scalar), both ops fix round-to-nearest and are marked
TVectorizable, so T.Parallel loops containing them keep their vectorized
loads/stores:
- scalar: __fmul_rn/__fmaf_rn (fp32), __dmul_rn/__fma_rn (fp64),
__hmul_rn/__hfma (fp16/bf16)
- vectorized on SM100+: packed x2 helpers tl::mul2/tl::fma2
- vectorized elsewhere: per-lane scalar calls
- CuTeDSL: tl.ieee_fmul/tl.ieee_fmaf with rounding="rn", expanded per
lane through a register tensor for vector dtypes
The packed x2 pair-decomposition logic is extracted from
PrintVecBinaryOp into a reusable EmitPackedX2Call helper (no behavior
change), alongside a new EmitPerLaneScalarCall fallback.
* [BugFix] Always vectorize `T.Parallel` loops
* Fix CI failures
* Restore always-vectorize accidentally reverted by the CI-fix commit
f447919b captured a transient working-tree state from a concurrent
baseline-vs-patched benchmark flip, reverting the gate removal in src/
while keeping the new-policy tests: CI then ran the old gated compiler
against them (scalar var-loop -> no float2; positional if/for check
broken by the copy loops the larger tile introduces).
Bring back the parameter-less LowerParallelLoop and rewrite the issue
1734 assertion as containment: the guarded loop must sit inside the
hoisted if block, which holds on both sm90 and sm100 regardless of how
many copy loops precede it.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Some ldmatrix lane maps step beyond the final SM75 MMA micro-tile, and
column overflow must carry into the next row. Layout remapping previously
linearized the source indices and optional offset before converting them
back to multidimensional coordinates, which normalized these addresses as
a side effect. Since #2836 preserves the multidimensional indices, the
out-of-range coordinates reach the swizzle layout unchanged and the native
SM75 MMA GEMM paths fault on illegal shared-memory accesses.
Wrap each complete row-major source offset within its operand region in
ldmatrix_a and ldmatrix_b before rebuilding the source coordinates.
Leading dimensions remain untouched, so the wrap cannot cross into another
stage of a multi-version buffer.
When the coordinates are provably in range, the simplifier removes the
wrap. SM70 and SM80 generated sources remain byte-identical, while the
SM75 addresses match the pre-#2836 address mapping.
Add multi-stage SM75 GEMM coverage and a lowered-offset bounds regression
that cross-compiles explicitly for sm_75 and therefore runs on any CUDA
runner.
_tir_packed_to_unsigned_convert_with_zeros performed the zero-point subtraction in the unsigned packed storage dtype before casting to the requested output dtype. Values below the zero-point therefore wrapped modulo the storage width.
Cast the unpacked field and zero-point to the output dtype before subtraction so asymmetric dequantization preserves negative values.
Add CUDA regression coverage for uint32-packed 4-bit values decoded with a non-zero zero-point.
Fixes#2947
* [CUDA] Unify TMA copy lowering on CuTe algebra and support general swizzled layouts
Every descriptor-based TMA path (bulk copy, atomic add, gather4/scatter4,
im2col) now shares one CuTe-algebra box decomposition. Boxes follow the
driver contract (inner box bytes at most the swizzle span, per-dim caps),
and the recovered swizzle is widened on demand under an equivalence proof,
so general annotated layouts -- including phase-shifted swizzles -- lower
correctly.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* Address review concerns
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* [Compiler][Z3] Bump 3rdparty/tvm to materialize Z3 solvers lazily
Z3Prover::Impl created a z3::solver eagerly in the Analyzer constructor
and translated every Bind/EnterConstraint into z3 constraints up front.
Analyzers are scratch objects created in huge numbers during lowering (a
reducer-heavy kernel compile creates ~67k of them, ~100k binds), while
only ~0.3% ever reach the Z3 fallback of CanProve, so the eager setup
dominated layout-inference-heavy compiles.
With the bump, binds and constraints are journaled and the solver is
created and replayed on first use. End-to-end compile of the worst
regressing reducer kernels drops 40-60%, and tilelang.lower roughly
halves across the board. Language suite 1168 passed / 137 skipped,
reducer_v2 75 passed, arith 89 passed; generated CUDA byte-identical
under disabled ASLR.
tvm changeset: https://github.com/tile-ai/tvm/pull/68
* [Cleanup] Drop Z3ContextScope wrappers
With the bumped tvm, every materialized Z3 solver owns a private context
(tile-ai/tvm#69), which subsumes the per-compilation isolation these
scopes provided; the scope functions are now documented no-ops. Remove
the wrappers around lowering and grouped compilation, and rewrite the
scope-lifetime arith test to assert what it actually guarded: a cloned
Analyzer's Z3 handles stay valid after the source is destroyed.
Validation (cold caches): layout suite 298 passed (including the
z3-context isolation tests, which cover the grouped-compile path),
language suite 1172 passed / 137 skipped, arith 89 passed, reducer_v2
75 passed.
* Point 3rdparty/tvm at the merged lazy-Z3 changes
tile-ai/tvm#68 and tile-ai/tvm#69 are both merged; replace the PR-branch
pin with the tilelang_main commit.
[Layout] Preserve vectorized reducer update plans
Treat reducer_update as a transparent expression during layout vector-shape analysis. This lets contribution loads determine the planned width instead of forcing update nests onto scalar-shaped layouts before reducer materialization.\n\nAdd a CUDA regression that checks the shared-to-fragment copy remains packed and verifies the reducer result.
math.h aliases several half-style names to cutlass::fast_* when a kernel
uses selected fast-math intrinsics. CUTLASS provides only float and double
overloads for fast_log, fast_sqrt, fast_sin, and fast_cos, and adds half_t
overloads only for fast_exp and fast_tanh. Overload resolution can therefore
leave a float result that cannot copy-initialize the CUTLASS wrapper.
Once a kernel includes math.h, for example through T.__exp, these aliases
also affect matching plain math calls in the same kernel. Their compilation
could therefore depend on whether another intrinsic pulled math.h in.
Add half_t and bfloat16_t overloads for fast_log, fast_sqrt, fast_sin, and
fast_cos, plus bfloat16_t for fast_tanh. Each evaluates in float32 and
converts back, following the existing fast_exp(bfloat16_t) bridge. Calls
that already compiled retain the same numerical path.
Add compile-and-run coverage for both wrapper types.
The probe inverts a candidate loop layout into symbolic per-thread
coordinates plus a symbolic thread variable, then checks whether any
fragment physical index still contains the thread variable. Proving a
strided distribution thread-free is range-dependent (e.g.
`(o0*512 + t*4 + o1) // 512 -> o0` only holds for t*4 + o1 < 512), so
with unbound ranges Simplify keeps the thread term and even an identity
access (loop layout == fragment layout) is falsely charged as a register
spill. The spill pricing then steers layout inference away from
conflict-free strided shared-memory loads toward chunk-contiguous
layouts whose 16-word stride bank-conflicts on every shared load,
observed as a 1.23-1.29x latency regression on a production
KL-divergence kernel (8.8ms vs 6.8ms at n=8192, d=262144 on H800).
Bind every per-thread output coordinate to its output-shape range and
the thread variable to the loop's thread range/extent so the legitimate
cancellations go through. Genuine thread dependence is unaffected: range
information only enables valid simplification, it cannot remove a real
thread term. Regression test distilled from the affected kernel.
* Steer reducer destinations during layout inference
* Remove extra reducer steering attempt filtering
* handle reducer problem3
* bug fix
* [Refactor] Give reducers a first-class PartialFragment layout solved by layout inference
The reducer's physical plan (narrow vs wide partials) used to be decided
after inference by ReducerPlanAndMaterialize, with dst-steering emulating
the missing propagation through side channels (wake tables, hint plumbing,
ownership gates); the same induced layout was computed twice and kept
consistent by convention only, and the reducer split the use-def graph
into disconnected components.
This makes the reducer's layout a first-class solver value that fully
determines the physical plan:
* New layout kind `tl.PartialFragment` (PartialFragmentNode : FragmentNode):
Fragment algebra whose replication coordinate carries an explicit
combine decomposition — under a low-bits convention enforced at
construction, `_rep % combine_size` enumerates addend lanes the finalize
collective must reduce and `_rep / combine_size` enumerates equal-value
copy groups. (storage algebra, combine width) is the complete plan: the
collective steps are derivable from the node alone (CombineSteps), the
FullParticipant object (combine == replicate == participants, thread ==
_rep) is the wide fallback, and combine width is part of layout identity
(same-storage different-decomposition proposals widen at the commit
point instead of unifying silently). A boundary-preserving condensation
(CondenseReplicateVarKeepingBoundary) keeps the convention true through
projection; kind-aware IsEqual in both dispatch directions prevents
silent unification with plain Fragments.
* Update nests (ParallelOp) record their reducer_update sites and, once
their loop layout solves, propose the induced partial (per-site
AnalyzeReducerUpdateSite proofs) or the participant-wide plan. The
commit point resolves multi-site disagreement with a monotone
widen-on-conflict rule; zero-update epochs are seeded wide before any
inference level runs; tl.reducer_force_baseline pins the proposal.
* finalize_reducer reads the solved partial and steers an unconstrained
destination to its post-collective reading. The wake table, the
LayoutInferArgs.reducer_update_sites hint plumbing and the
reducer_use_list_ component glue are deleted: reducers travel ordinary
use_list_ buffer edges.
* ReducerPlanAndMaterialize reads the plan off the node: storage =
AsPostCollective(), steps = CombineSteps(), wide lowering iff the node
is the FullParticipant object. TryNarrowPlan is retired; what remains
(DeriveNarrowLowering) derives only the lowering details that live
outside the layout — the packed-lane heuristic and the
destination-containment gate (which still demotes gracefully to the
wide plan, preserving the unsatisfiable-steering fallback contract).
The destination override chains are kept as a verified-no-op fallback
with retirement telemetry.
* T.annotate_layout accepts a PartialFragment on reducer buffers, pinning
plans the automatic solution cannot reach (e.g. a middle plan between
narrow and wide). Annotated partials are strict: a disagreeing update
site raises with the narrow-rejection reason, and
ComputeReducerPinnedCandidate back-propagates the annotation through
the update loop into its input fragments.
* Layout cost models price register-array spills: a fragment access whose
physical slot index depends on the thread variable (detected by the
shared FragmentThreadIndexProbe, extracted from
HasThreadDependentFragmentIndex) demotes the whole array to local
memory, charged as byte-denominated traffic in both models' `mem`
channel so it competes with — rather than vetoes — global-bandwidth
estimates. ChooseBestCandidate's anti-spill preference now applies
regardless of where the loop stores (a fragment->fragment publication
copy spills exactly like a fragment->global one). With generation and
pricing combined, inference produces the wide publication chain's
replicated layouts on its own: the reducer_v2 suite and the
mhc_post_bwd production kernels pass with the destination override
fully disabled.
Testing (all with the kernel cache disabled): reducer_v2 suite 70 passed
(new coverage: pinned middle plans, annotation steering input fragments,
conflict diagnostics, combine-step roundtrip, same-storage
different-combine widening at commit); language suite 1261 passed;
layout-inference cost-model tests and maint/layout_inference goldens
unchanged; generated CUDA for existing reducer kernels is byte-identical;
mhc_post_bwd correctness passes and its six benchmark sizes show
no latency change with zero STL/LDL in SASS.
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[Bugfix] Make TMA linear layouts region-aware to keep versioned slices contiguous
The inferred "linear" shared layout hoisted every 256-split quotient above
all dims, so a pipeline-versioned slice became two gapped chunks: LowerBulk's
bijection check fired on warp-specialized wide reductions, and the TMA
atomic-add path (which has no such check) silently read the other version.
MakeTmaLinearLayout now takes the copied region and orders modes as
[region-fixed dims, box quotients, box contents], so a fixed slice owns one
contiguous run of whole TMA boxes and the existing rest-loop machinery emits
one instruction per box over the true (unaligned) global shape. The
atomic-add path shares the helper and computes the region base offset through
the layout instead of hand-rolled row-major strides over the physical shape.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
A synchronous T.gemm(mbar=...) inside T.Pipelined waited on the wrong
mbarrier parity. InjectSoftwarePipeline annotated every copy/im2col/gemm
call with the ring parity (k / num_stages) % 2, but that law belongs to
the compiler-generated pipeline_mbar ring: an explicit mbarrier whose
arrive happens inside a tile op is deliberately excluded from
ExpandPipelineBarriers, so it remains a single slot that flips once per
iteration. The auto-emitted wait asked for a parity that had already
completed, returned immediately, and the next iteration's copies
overwrote a shared-memory stage the still-running MMA was reading
(wrong results observed at num_stages >= 3).
Annotate only the ops that consume the phase -- those that auto-emit a
paired arrive+wait on an explicit mbarrier: sync gemm with an mbar, and
copies carrying an explicit "barrier" that are not is_tma_copy -- and
annotate them with the logical iteration parity (k - min) % 2, written
against the original loop var. EmitImpl's existing per-block
Substitute(loop_var -> normalized_access_index) specializes the
expression for every prologue/body/epilogue clone, since the tirx
mutators already rewrite PrimExprs inside call annotations.
Pipeline-managed TMA copies keep their ring parity in the explicit
waits the pass inserts, and the explicit-async T.tcgen05_gemm never
auto-waits, so neither is annotated.
Drop the matching ring branch from LowerTileOp's fallback: it keyed on
the "num_stages" annotation, which InjectSoftwarePipeline strips, so it
was dead on pipelined loops and wrong on unpipelined ones -- a
compiler-generated per-site barrier flips once per loop iteration
regardless of pipeline depth. The fallback is now the zero-based epoch
of the nearest enclosing serial loop.
Add transform-level tests asserting the annotated phase and the emitted
wait parities on sm_90a, so the fix is covered without sm_100 hardware;
the loops start at a nonzero min to pin the epoch normalization. Add a
tcgen05 correctness test running the failing kernel at num_stages 2
and 3.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Remove an unconsumed cluster-planning pass, duplicate sparse GEMM and fast-math implementations, the unused JIT environment module, and the superseded autotuner timeout machinery. Document and test the surviving portable timeout behavior.
TVM lowers 16-bit math operations to CUDA's half-style names, but CUDA
provides overloads only for native __half and __nv_bfloat16 values.
Calls with CUTLASS half_t and bfloat16_t wrappers therefore failed to
compile.
Add explicit forwarding overloads for hexp, hexp2, hexp10, hlog, hlog2,
hlog10, hsin, hcos, hfloor, hceil, hrint, and htrunc. Each overload
converts the wrapper to its native CUDA type, calls the existing
intrinsic, and wraps the result back in the original type. Fold the
previous standalone bfloat16 hexp bridge into the complete overload
set. Codegen and the fast-math path remain unchanged.
Add scalar runtime coverage for both wrapper types using materialized
temporaries, separate valid input domains, generated-call assertions,
and torch references.
* feat(tools): add official no-GPU compile-only entry (#2913)
TileLang has no first-class compile-only path for CPU-only CI.
This adds the official entry so later CE work can emit inspectable
kernel source without a GPU.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* feat(tools): strengthen optional compile-only CUDA target (#2913)
`--target cuda` should emit real CUDA C when the wheel has codegen FFI,
and soft-fail clearly on Metal-style wheels instead of an AttributeError.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* fix(tools): honor module order and CUDA target options (#2913)
discover_prim_func skipped an earlier PrimFunc when a later @jit existed,
and option-bearing CUDA strings skipped the FFI soft-fail.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
* fix(tools): drop stale --output_file when compile-only fails (#2913)
CE reuses the same output path; a failed run must not leave yesterday's
assembly. Unlink before compile, matching the CuTe wrapper.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
* fix(tools): reject JSON auto and pin bare CUDA JSON (#2913)
CodeRabbit: '{"kind":"auto"}' bypassed the auto ban, and
'{"kind":"cuda"}' skipped the sm_80 pin. Apply the same rules as
string targets.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
* fix(tools): refuse --output_file when it is the input (#2913)
Unlink-before-compile deleted example.py if someone passed the same
path twice. Compare resolved/samefile paths first.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* fix(tools): emit #line from compile-only when the pass exists (#2913)
Pass TL_EMIT_LINE_DIRECTIVES so CE can map generated C back to the
example. Older wheels do not register the key; skip it there.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* test(tools): require #line to name the input file and store line (#2913)
CodeRabbit: a bare "#line" token can pass on an unrelated directive.
Parse `#line N "file"` and match the example path plus the marked store.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* test(tools): give in-process #line lookup its own marker (#2913)
_marker_line scans the test file, so a shared line_marker_store hit the
example string (line 37) instead of _add (line 50). ROCm CI asserted that.
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
* docs(tools): document the compile-only CLI (#2913)
Co-authored-by: Oliver Zhang <LibertychaserUS@users.noreply.github.com>
* docs(tools): document remaining compile-only helpers (#2913)
Co-authored-by: Oliver Zhang <LibertychaserUS@users.noreply.github.com>
---------
Signed-off-by: Oliver Zhang <zhiyuanzhang0904@163.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Oliver Zhang <LibertychaserUS@users.noreply.github.com>
T.shfl_sync, T.shfl_xor, T.shfl_down, and T.shfl_up lower to CUDA's raw
__shfl_*_sync builtins. CUDA has no exact overloads for the narrow
floating-point wrappers TileLang emits, so calls either resolved to the
float overload and returned a float that could not copy-initialize the
wrapper, or were ambiguous. Either way the shuffle failed to compile.
Add exact non-template overloads for half_t, bfloat16_t, float_e4m3_t, and
float_e5m2_t. The overloads carry the raw bits through CUDA's uint32_t
shuffle and preserve mask, lane, delta, and width semantics. A local macro
generates the four builtin overloads per type and is undefined immediately
after use, leaving codegen and native CUDA overload resolution unchanged.
Add exact lane-permutation coverage for all four shuffle operations and all
four wrapper types in both materialized-temporary and inline forms.
Fixes#2569
A tile-region `T.atomic_add(..., memory_order=...)` lost its ordering
whenever the add auto-vectorized: the vectorizer rebuilt the wide op from
only `{dst, src}`, so `AtomicAddx2`/`AtomicAddx4` were emitted without the
order argument and fell back to their `relaxed` default. The identical call
honored the order when it lowered scalar-wise, so whether the requested
ordering survived depended only on whether the per-thread extent was
divisible by the atomic vector width.
Carry the trailing operands through the wide reconstruction. Codegen and
the device helpers already accept the order (`args.size() > 2` guard, and
`int memory_order` parameters that branch to ordered PTX), and
`atomic_addx2/x4_elem_op` are declared with `num_inputs = 3` -- the
vectorizer was simply under-filling the op. The order is copied as-is
rather than visited, since it is a lane-invariant scalar that must not be
broadcast across the vector lanes.
Carrying the order also wakes up inline PTX that was previously
unreachable: with the order always defaulted to relaxed, the non-relaxed
branches were constant-folded away before ptxas ever saw them. Three
problems surface once they are live, all verified with ptxas 13.3:
1. `atom` with vector operands requires sm_90, so the ordered fp16 pair
failed to assemble on sm_80/sm_86/sm_89 -- and fp16 vectorizes to x2 on
every target, so that is every ordered fp16 region atomic_add. Use the
packed `f16x2` form instead of `.v2.f16`: it encodes the same pairwise
add but is available from sm_60, so the wide op stays a single atomic
instruction on every supported target rather than needing an arch split.
sm_89 SASS confirms a native ATOMG.E.ADD.F16x2.RN.STRONG.GPU. This also
drops the per-half pack/unpack, since a half2 reinterprets directly to
the 32-bit operand.
2. The bf16 v2 asm was missing the `.noftz` modifier that ptxas requires
for `.bf16`, so `atom.{release,acquire,acq_rel}.gpu.global.add.v2.bf16`
failed to assemble on every arch, sm_90 included. The scalar bf16
helper next to it already spells it correctly.
3. bf16 has no packed equivalent below sm_90 -- `bf16x2` and `.v2.bf16`
both require sm_90, as does scalar `atom.add.bf16` -- so unlike fp16 it
genuinely cannot express the ordering in one instruction there. Guard
the sm_90 path and emulate the ordering below it with device-scope
fences around a relaxed bf16x2 atomicAdd. `__threadfence()` lowers to
`fence.sc.gpu`, the same device scope as the `.gpu` qualifier on the
ordered PTX, and is two-way where release and acquire need only one-way,
so this is conservative rather than lossy. `AtomicAddx2Ret` takes the
same guard, since it shares the helper. Nothing reachable instantiates
the bf16 Ret overloads with a non-relaxed order today, so that path has
no test; it is guarded to keep the file consistent and to stay correct
once region-form `return_prev=True` is implemented.
Add a regression test in fp16 and bf16. It disables the kernel cache,
since the key does not cover the native library by default and would
otherwise serve stale pre-fix source, and it runs the kernel so the
ordered wide path is checked for correctness and not just for the emitted
argument. The bf16 case covers the pre-sm_90 fenced branch on Ampere/Ada
and the sm_90 instruction on Hopper.
Also correct the stale memory_order sentence in atomic_max and
atomic_min, which forward the order through the shared region lowering
just like atomic_add does.
Fixes#2688
Signed-off-by: Srijan Keshri <srijankeshri007@gmail.com>
[CUDA] Add FP8 min/max overloads
CUDA codegen emits unqualified min/max calls for TIR MinNode and MaxNode
expressions. FP8 operands promote to float, but the FP8 wrappers have no
implicit conversion back from the result, so T.max, T.min, and T.clamp
failed during nvcc compilation.
Add tl::min/max overloads for float_e4m3_t and float_e5m2_t. The overloads
compute through fminf/fmaxf and convert the result back to FP8. ADL lets the
existing scalar and per-lane vector codegen resolve these overloads without
FP8-specific codegen branches.
Add runtime coverage for both FP8 formats and compile coverage for vectorized
expressions and the e4m3fn spelling.
Fixes#2985
