* [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
T.__ldg failed to compile for float16 and bfloat16 buffers because CUDA
provides __ldg overloads for __half and __nv_bfloat16, while TileLang uses
CUTLASS half_t and bfloat16_t for these scalar types.
Add forwarding overloads in common.h that reinterpret the bit-compatible
pointers as the native CUDA types and wrap the loaded values back in the
CUTLASS types. The native 16-bit __ldg is available on SM60+, so no
architecture-specific fallback is needed.
Update the stale 32-bit-only API wording and add exact round-trip coverage
for both dtypes.
* [Op] Add read-before-write semantics to tile operators
GetAccessRegions() answers a may-read question: could this operator read this
region? It is deliberately conservative, which is correct for the dependency
analysis it was built for. There is no way to ask the stricter question of
whether an operator consumes a region.s preexisting contents before
establishing a value of its own.
GetReadBeforeWriteRegions() asks it. The default returns GetAccessRegions()
.reads, so operators whose read set does not depend on an argument value need
no change, and existing consumers of GetAccessRegions() are untouched.
Gemm and GemmSP override it. GetAccessRegions() lists the accumulator as a
read whenever the clear cannot be proven (!is_one(clearAccum_)), which is
right for pipelining: a conditional clear genuinely is a read dependency. It
is wrong for a read-before-write question, because the common pipelined idiom
clear_accum=(k == 0) is neither literal true nor literal false. The overrides
use is_zero() instead.
Note that operators such as ReduceOp need no override: its clear is a plain
bool rather than a PrimExpr, so the default is already exact.
* [Transform] Add VerifyBufferInit, a general buffer-initialization check
Reading a non-global buffer that nothing has written yields whatever those
locations last held. That is undefined behaviour: it can look correct on one
architecture and produce NaN on another (issue #2936).
VerifyBufferInit walks a PrimFunc before LowerTileOp, where tile operators
still declare their access regions, so writes performed by T.copy, T.clear and
friends are visible rather than opaque. It warns and returns the function
unchanged; it never fails a build.
The analysis is conservative by construction, in three ways:
- A call the pass cannot interpret may write through its arguments, so it is
assumed to. TCallEffectKind distinguishes these from pure calls; an
operator with no registered effect kind is assumed to write. This also
covers tl.access_ptr, which keeps a BufferLoad rather than a data Var.
- Shared memory ignores source order. A warp-specialized kernel fills a shared
buffer from a producer branch that appears after the consumer branch that
reads it; the two run concurrently and coordinate through barriers, so
source order is not an approximation of happens-before. For shared scopes
the question is whether any *other* operation writes the buffer at all.
Per-thread storage keeps order-sensitive tracking, where source order is
that thread.s execution order.
- Global memory and mbarrier scopes are never reported.
Excluding the reading operator.s own write is what keeps #2936 reported: a
gemm.s only writer of its accumulator is itself.
* [Transform] Enable the buffer-initialization check in every pipeline
Wired into all five backends immediately after LayoutReducer, and on by
default with tl.disable_buffer_init_check to opt out.
The slot matters in four ways. LegalizeNegativeIndex has run, so buffer
indices are canonical. LowerTileOp has not, so tile operators still carry
their access regions. InjectSoftwarePipeline has not, so loop bodies are in
source order rather than rotated across stages. And LayoutReducer has run:
before it, a tl.tileop.finalize_reducer call carries a single argument while
its builder reads args[1], so ParseOperator throws on it. Every pre-existing
ParseOperator caller sits after that point, which is why the incomplete form
had not been parsed before.
22 tests cover the gemm cases, the general cases, the two scope-ordering
rules, report aggregation, and the config plumbing end to end.
* [Transform] Route both buffer-init walks through one opaque-write rule
The whole-function collector and the ordered walk each decided for
themselves what an uninterpretable call may write, and they disagreed.
The collector recorded any BufferLoad argument, so a read-only
tl.access_ptr counted as a writer and could suppress a real finding on
shared memory, while a write through tvm_access_ptr passed a data Var
rather than a BufferLoad and was missed entirely. It also applied the
rule to pure calls, whose arguments only read.
Both now share ForEachOpaqueWrite and the same escape predicate, so the
access-mask and pointer-form handling cannot drift apart again.
Also adds a regression test for the unguarded loop-carried read. That
shape is reported and the report is correct: the first iteration reads
the fragment before the write at the bottom of the body reaches it. The
case the analysis does not decide is a read the loop guards out of its
own first iteration, which is path sensitivity rather than ordering.
Verified: 23 verifier tests pass; testing/python/language is 997 passed
with the same 5 pre-existing failures and zero warnings; all 284 example
scripts emit one warning, the known uninitialized lse_max_local in
examples/flash_decoding/example_mha_inference.py.
* [Transform] Record a store's destination after evaluating it
A BufferStore reads its own right-hand side before it establishes
anything, so marking the destination written first hid every
self-referential read: `x[i] = x[i] + 1` with nothing written earlier
went unreported.
I had defended that ordering on the grounds that flagging
`idx = T.if_then_else(cond, i, idx)` turns the check into
definite-assignment analysis over conventional idioms. Looking at the
two examples that motivated it, that defence does not survive.
examples/grouped_gemm/example_grouped_gemm_fwd_ptr.py writes
`cur_batch_idx = 0` before exactly this loop; example_grouped_gemm_fwd.py
and example_grouped_gemm_bwd.py do not. The initializer is already there
in the sibling, so this is an oversight in two of three copies rather
than an idiom the check misreads.
Cross-thread scopes are unaffected: they ignore source order and ask
only whether another node writes the buffer.
Verified: 24 verifier tests pass; testing/python/language is 997 passed
with the same 5 pre-existing failures and zero warnings; all 284 example
scripts yield 3 warnings in 2 places, the lse_max_local reduce
destination and the two cur_batch_idx sites, with no false positives.
Note the two grouped_gemm examples also fail their own torch.allclose
check, but that is unrelated: initializing cur_batch_idx leaves the
mismatch at an unchanged 0.25 max absolute error over ~3.4k of 1572864
elements.
* [Transform] Discount a store's own write when checking its reads
CheckRead already took the reading node, so a cross-thread scope could
ask whether some other node writes the buffer. That is the rule keeping
a gemm accumulating into an untouched accumulator reported. Loads passed
no reader, so a store that read its own destination counted itself as
its own initializer: `s[i] = s[i] + A[i]` on shared memory was silent
while the identical code on a fragment was reported.
The enclosing store is now published while its subexpressions are
visited. A store is recorded as a writer of nothing but its own
destination -- an opaque call inside the value is recorded under the
call node -- so discounting it cannot suppress a report for any other
buffer.
Also covers the reducer v2 idiom introduced by #2940. reducer_init,
reducer_update and finalize_reducer inherit the default
GetReadBeforeWriteRegions and need no override, which is the same
result the pass already gets for T.reduce.
* [Transform] Exempt a parameter buffer from the check at any scope
SeedParams treats every parameter as written by the caller, but the
cross-thread branch of CheckRead asks only whether another node in the
body writes the buffer, and the caller is not one. A shared-scope
parameter therefore reported while an identical local-scope parameter
stayed silent.
Record the parameters in their own set and admit them ahead of the
scope-dependent question, so WrittenByAnotherNode is left untouched for
every buffer the body has to establish itself.
* [Transform] Print the buffer name and widen the parse fallback
Streaming the Buffer sent a one-word diagnostic through the reflection
repr printer. The name is a String with a direct stream overload and is
what the message wanted either way, so print it directly.
TryParseOperator caught tvm::ffi::Error, which every builder failure
seen so far raises. It derives from std::exception, so catching the base
extends the documented degrade-rather-than-abort rule to the failures
that would not.
* [Example] Initialize the buffers the new check reports
Three sites, all read before anything writes them. None is new and none
is known to misbehave; they are benign on the inputs these examples use,
which is why nothing has surfaced them.
grouped_gemm fwd and bwd allocate `cur_batch_idx` and immediately carry
it forward through `T.if_then_else(cond, i, cur_batch_idx)`, so a false
predicate on the first iteration returns whatever the variable held.
`cur_batch_idx = 0` is what example_grouped_gemm_fwd_ptr.py already does,
and has done since its first commit (#1923) — a second author writing the
same loop initialized it unprompted.
flash_decoding passes an uninitialized `lse_max_local` to
`T.reduce_max(..., clear=False)`, which makes the destination a read by
the operator's own semantics. It is the only write to that buffer and it
sits two lines below `T.clear(lse_logsum_local)`, so `clear=True` is what
was meant.
Verified: all 284 example scripts now compile with zero warnings.
example_mha_inference.py still reports "All checks passed!" and its
output is bit-identical to the current form on an RTX 3060 Ti.
The two grouped_gemm examples still fail their own torch.allclose check.
That is a separate, pre-existing issue untouched here: the mismatch is
unchanged at 0.25 max absolute error over ~3.4k of 1572864 elements
whether or not cur_batch_idx is initialized.
* [Refactor] Read the buffer-init config inside the pass
The pass function already receives the PassContext, so it can gate
itself on tl.disable_buffer_init_check rather than have every backend
pipeline consult a Python helper first. This matches ThreadSync, which
reads tl.disable_thread_storage_sync the same way and is scheduled
unconditionally.
Drops should_enable_buffer_init_check and the five gated call sites.
The two tests that asserted on the helper now call the pass directly,
which also covers the case the Python gate never did: invoking
VerifyBufferInit outside the pipeline under a disabling PassContext.
* [Lang] Reducer v2: first-class deferred reduction epochs with planned physical lowering
Fixes#2408 at its root. v1 implemented reducers as fragment buffers with
a special layout, so no pass owned the contribution-multiplicity
semantics: loop-layout replication silently multiplied contributions
(128 threads summing 8 elements returned 576 instead of 36), the
finalize width had to be guessed from the storage layout's
ReplicateExtent, and every compiler stage carried buffer-identity
special cases to compensate (see #2897, #2881).
Reducer v2 makes the reducer a first-class deferred reduction epoch.
Core invariant: `T.reducer_update` is the single source of contribution
multiplicity; physical layout can never change how many times a logical
contribution is combined.
New IR contract (`local.reducer` virtual scope + three first-class ops):
acc = T.alloc_reducer(shape, dtype, op="sum", seed=None)
T.reducer_init(acc)
T.reducer_update(acc[indices], contribution) # inside T.Parallel
T.finalize_reducer(acc, dst) # out-of-place
Pipeline:
* VerifyReducerEpoch (early): lifecycle state machine + opaque-access
rules; ordinary reads/writes, clear/fill, aliasing, updates outside
T.Parallel, double init/finalize are compile errors.
* ReducerPlanAndMaterialize (after LayoutInference; loop layouts are
read-only inputs): decides storage, communication, and execution
multiplicity per epoch, then materializes the first-class ops into
ordinary IR. Two physical plans:
- Narrow plan: project the reduction axes (parallel vars absent from
the update indices, plus loop replication) out of the update site's
loop layout to obtain compact per-thread partial storage; updates
run unguarded on every replica; the collective covers only the
thread-expression splits sourced from reduction axes (reusing the
ComputeReducerLayout / CollectThreadReduceSteps machinery T.reduce
is built on). Zero splits means zero collectives (LocalComplete).
Proof obligations: direct var-to-dim indices, replica-safe
contribution values, full-block coverage, a single power-of-two
step, structurally equal plans across sites, destination-layout
containment (with a controlled override for provably unconstrained
single-slot destinations).
- Wide plan (FullParticipant baseline, always available): one full
logical-shape partial per participant, canonical-replica-guarded
updates, participant-wide AllReduce, seed combined exactly once
post-collective. Any narrow-plan proof failure falls back here
deterministically - never to a compile error or wrong code.
* Generic multiplicity marker `tl.parallel_multiplicity`: PartitionLoop
lowers it to a REP == 0 guard knowing nothing about reducers; the
fully_replicated_reducer_buffers compatibility hook is retired.
* FinalizeReducerOp takes explicit plan operands (reducing_threads,
scale), closing the derive-width-from-ReplicateExtent channel.
* VerifyReducerConsumed (after LowerTileOp): surviving v2 constructs
are compile errors, never silently-wrong code.
* Pass config `tl.reducer_force_baseline` forces the wide plan for
differential testing.
Legacy compatibility and v1 removal: CanonicalizeLegacyReducer rewrites
v1 syntax (clear + read-modify-write stores + in-place finalize) into
v2 ops on a strict whitelist; non-identity fills on idempotent combines
are forwarded as one-time seeds, non-zero sum fills are rejected, batch
annotations are forwarded, and post-finalize reads are redirected to a
synthesized destination. With the shim in place, the entire v1 lowering
path is deleted: LayoutReducer, the PartitionLoop/LowerParallelLoop
reducer-buffer plumbing, the ParallelOp rep=ALL special cases, and the
vectorizer's reducer parameter chain.
Testing (H100, kernel cache disabled): 128 passed / 2 skipped across
the legacy reduce suite (93, unchanged, through the shim), the new
reducer v2 suite (22, incl. the 36-vs-576 regression, narrow-plan
codegen assertions, lifecycle diagnostics, forced-baseline
differentials), and the copy suite; the pipelined GEMV example (legacy
syntax) remains numerically correct.
* [Refactor] Consolidate reducer ops into src/op/reducer.{h,cc}
FinalizeReducerOp is no longer a standalone legacy op — it is the
collective that ReducerPlanAndMaterialize emits — so it moves next to
the first-class reducer ops. Its combine-op field now uses the same
ReducerV2OpType enum as the planner, removing the implicit
static_cast coupling between two parallel enums.
The former LayoutReducer pass files are gone entirely: the
ReducerInfo/ReducerRepType annotation types had no consumers left, and
the legacy `reducer_info` annotation key (still read by the
canonicalization shim and the data-race verifier) moves to
op/reducer.h. Stale includes in layout_inference / lower_tile_op are
dropped.
No functional change; 115 reducer tests pass.
* [Doc] Clarify the intent of tl.reducer_force_baseline
The switch is not a workaround for expected narrow-plan bugs; it exposes
the reducer design's canonical reference lowering for differential
testing, field escape-hatch use, and plan-choice A/B measurement.
* [Layout] Opaque tile-op loops must not self-plan their layout
A parallel loop whose body contains an opaque tile-op call (e.g. a
reducer update) has no visible write to anchor layout inference. In
free mode such a loop used to self-plan a partition, and the
register-count-based attempt selection then preferred it: the plan
minimizes this loop's own footprint (e.g. one x-vector element per
thread) while silently forcing a column-major layout onto the fragment
operands it reads, degrading their producing copies from vectorized to
scalar accesses. On a 4096x4096 fp16 GEMV this cost 40% end-to-end
versus the legacy reducer lowering.
Fix (generic, not reducer-aware): when a loop carries an opaque
tile-op call, prefer propagating its layout from a fragment operand
that is addressed purely by the parallel loop vars, deferring while
such an operand is still unsolved so the producing ops fix the
operands first. Self-planning remains the fallback when no operand can
drive the loop (e.g. operands indexed by inner serial reduction vars,
which require the loop to own the operand layout instead) or when
operand propagation cannot express the loop's iteration space.
With the fix the GEMV benchmark matches the legacy lowering (0.075 ms
vs 0.074 ms) and narrow-plan row reductions stay ahead (AllReduce<16>
with one partial slot per thread, ~6% faster than legacy); 129 tests
pass across the reducer v1/v2 and copy suites.
* [Lang] Reducer v2 stage 4: narrow-plan coverage, packed accumulation, bitwise ops
Narrow-plan coverage and observability:
- tl.enable_reducer_plan_verbose pass config: INFO-level per-epoch plan
decisions (narrow/wide, collective widths, rejection reasons).
- Multi-step collectives: interleaved thread splits lower to one AllReduce
per step instead of falling back to the wide plan.
- Destination override generalized: multi-slot destinations and transitive
copy chains (e.g. fp32->fp16 staging hops) adopt the induced layout when
the whole chain is unconstrained; ParallelOp candidate selection falls
back to the operand-compatible layout when the register-preferred
candidate conflicts with an overridden buffer.
- Update-index generalization: permuted loop-var indices and constant unit
dims (acc[j, i], acc[0, i]) rebuild the induced fragment over the
reducer's own dim order; all-constant scalar indices keep the synthetic
unit input ComputeReducerLayout leaves behind.
- Seed and batch now work under narrow plans: the finalize lowering
combines the seed once per slot after the collective, and explicit-width
steps compose with batched AllReduce.
Packed partial accumulation (16-bit floats):
- Single-site narrow epochs on fp16/bf16 split each per-thread partial
into two lanes keyed by the parity of an on-thread reduction source (the
innermost enclosing serial loop, or the innermost parallel reduction dim
when its low bit stays on one thread), halving the serial combine
dependence chain.
- A per-thread fold recombines the lanes into the plain induced storage
before the (unchanged) collective, so the plan's communication and
destination proofs are untouched. Lane parity never changes which thread
a contribution lands on, so the plan stays correct under any layout.
Bitwise combine ops:
- alloc_reducer(op="bitand"/"bitor"/"bitxor") for integer dtypes under
both plans (identities: all-ones / 0 / 0). Float dtypes are rejected at
allocation time; the legacy v1 path stays sum/max/min only.
Tests: reducer v2 suite grows to 38 (packed structural + numeric +
differential, bitwise under both plans, permuted/unit-dim/scalar indices,
narrow seed); the v1 codegen batch tests pin tl.reducer_force_baseline
since narrow plans legitimately skip run_batch. v1 suite stays green (93).
* [Lang] Move the reducer seed to T.reducer_init(acc, init)
The logical starting value used to be declared at allocation time
(alloc_reducer(seed=...)) and smuggled to the planner through the
reducer_info_v2 block annotation. Declaring it where the epoch opens
reads naturally ("the reduction starts from this value") and makes it a
first-class operand of the init op:
acc = T.alloc_reducer((M,), dtype, op="sum")
T.reducer_init(acc, 100.0) # optional; default = combine identity
- ReducerInitOp accepts an optional second argument carrying the value;
the annotation now carries only the combine op.
- VerifyReducerEpoch checks the value's dtype at the init site (Python
numbers are auto-converted by the frontend).
- CanonicalizeLegacyReducer forwards v1 non-identity idempotent fills as
the synthesized reducer_init's argument instead of via the annotation.
- Semantics unchanged: physical partials always initialize to the
identity; the value is combined exactly once per logical output at
finalize time, so physical replication can never multiply it.
Tests: seed tests migrated to the new API plus a dtype-mismatch
diagnostic (v2 suite now 39); v1 suite stays green (93).
* [Layout] Name the operand-driven-loop contract: kTLPerIterationOp
The opaque-loop rule (6749631b) worked but read poorly: detection
sniffed the tile-op registry (minus a tl.region special case), the
policy lived in comments spread over three code blocks, and "can this
operand drive the loop?" was answered twice with different machinery
(a try/catch around the generic source-buffer choice, plus a separate
drivability scan for the defer decision).
Restructure without changing the architecture:
- New op attribute kTLPerIterationOp declares the contract at the op
definition site: the op executes once per iteration inside T.Parallel
and its region-mediated effects are invisible to statement-level
scans. tl.reducer_update is its only member today. Both consumers now
read the attribute instead of keeping their own lists: ParallelOp's
detection (previously registry sniffing) and the nested-loop
checker's _PARALLEL_SAFE_TILE_OPS whitelist (deleted).
- Free-mode inference for such loops is now a single decision:
FindOperandDriver returns the best solved drivable operand (indices
pure in the loop vars, most index dims first) or reports a pending
one; the loop adopts, defers, or falls through to self-planning. The
exception-as-control-flow try/catch is gone from the policy path (a
small guard remains for non-affine driver indices, which drivability
does not test). One behavioral refinement: when the generic
most-indices read buffer is not drivable but another solved drivable
operand exists, the loop now adopts that operand instead of falling
back to self-planning.
Verified: reducer v2 (39) + v1 (93) + copy suites pass; the reducer
benchmark (rowsum fp32/fp16, GEMV fp16) reproduces identical plans and
timings to the previous commit.
* [Lang] reducer_update is a per-iteration builtin intrinsic, not a tile op
reducer_update never was a tile op in substance: it executes once per
iteration INSIDE T.Parallel (no other tile op does), owns no layout (the
enclosing loop and the planner decide physics), and never lowers on its
own. Registering it as one forced two categories of friction: consumers
that had to exempt it (the nested-loop checker whitelist, ParallelOp's
registry sniffing before kTLPerIterationOp), and an argument convention
(tl.region wrapping) whose only real payload was the multi-dim indices.
Re-register it as the builtin `tl.reducer_update`:
- args are (acc[indices], value) with a plain BufferLoad target — an
update descriptor keeping structured indices for the planner. Every
access scan now sees the acc read for free; the write stays
undeclared, which is sound for reducers specifically: updates commute
(combine ops are associative/commutative) and VerifyReducerEpoch pins
init/finalize to straight-line code, so no cross-statement write
ordering exists to lose inside pipelined/parallel regions.
- ParseReducerUpdate replaces the TileOperator ctor; verifier, planner
and the legacy shim consume it. The verifier visits the target's
indices but not the target load itself (a reducer load anywhere else
is still the usual compile error).
- Pipeline statement classification counts kTLPerIterationOp calls as
compute (BufferRegionCollector, IsPureCopyStmt/GetSinglePureCopyTileOp
would otherwise classify an update-only loop as copy-like now that the
call no longer parses as a tile op); the shared IsPerIterationOpCall
helper also replaces ParallelOp's inline attribute lookup.
- The nested-loop checker exemption is gone entirely: a builtin has no
TLOpBuilder, so it passes the tile-op ban by construction.
Verified: reducer v2 (39) + v1 (93) + copy suites pass; the reducer
benchmark (all cases T.Pipelined, exercising the pipeline dependency and
classification paths) reproduces identical plans and timings.
* [Cleanup] Drop the TLPerIterationOp attribute for a direct predicate
The attribute named a category that does not carve reality at a joint:
nearly every builtin intrinsic "executes per iteration" inside
T.Parallel, so the name suggested a structural property while actually
being an opt-in marker with exactly one member. What its two consumers
really key on is reducer_update itself — the hidden write to a
deferred-layout buffer (ParallelOp inference) and compute-vs-copy
statement classification (pipeline planning).
Replace it with IsReducerUpdateCall next to the op declaration; both
consumers and ParallelOpNode::has_reducer_update_ now say plainly what
they check. A future sibling op (e.g. a rescale intrinsic) extends the
one predicate.
* [Doc] Record the operand-driven defer protocol's known limitation
Deferring assumes the pending drivable operand will be solved by its
producer. An operand with no producer at all (an update reading a
fragment nothing ever writes) defers on every free-level round and
surfaces as a compile-time "layout can not be inferred" ICHECK — never
wrong code, and only on degenerate input. A sound fix needs the
inference engine to order solving (or guarantee a final round); that
redesign is deliberately deferred to a follow-up rather than patched
into the protocol here.
* [Layout] Drop the operand-driven defer protocol from ParallelOp
The three-piece workaround (has_reducer_update_ detection, the
FindOperandDriver three-state ladder, free-mode defer) patched solve
ordering for anchor-less loops from inside a single op. The real defect
is the engine's register-count attempt scoring: the bias row-broadcast
counterexample reproduces the same mis-selection with no anchor-less
loop involved, so the protocol treated a symptom. Direction for the
proper fix (IO-aware cost model + declared solve ordering) is recorded
in the layout RFC.
Reducer-update loops return to the generic two-candidate free
inference. The candidate-validation fallback still bounds the damage:
fp16 GEMV 0.0615 ms vs 0.0736 ms v1 baseline (down from 0.0346 ms with
the protocol), rowsum unchanged across all sizes. reducer v2 (39),
v1 reduce (93), atomic, and copy suites pass.
* [Cleanup] Unwrap the free-inference branch's vacuous layout guard
The inner !loop_layout_.defined() check was the seam where the deleted
operand-driven block could adopt a driver layout before the generic
two-candidate path ran. The enclosing branch already requires an
undefined layout and nothing assigns one in between, so the guard is
always true now.
* [Cleanup] Drop the unreachable post-selection candidate fallback
ChooseBestCandidate already validates both candidates with the same
flags and returns a valid one whenever either validates (single-valid
returns it directly, both-valid picks among valid, both-invalid leaves
nothing to fall back to). The post-hoc re-validation could therefore
never flip the choice; an instrumented run across the reducer suites
and the rowsum/GEMV benchmarks confirms it never fires. It was added
alongside the destination-chain override work and mistakenly credited
for making multi-slot overrides safe — the chain registration is what
did that. Comments referencing the fallback updated.
* [Cleanup] Tighten loop_partition after the multiplicity-marker rework
- MultiplicityMarkerLowerer only rewrites statement-level AttrStmt
markers; derive from StmtMutator instead of StmtExprMutator so
expression subtrees are not traversed.
- Drop a duplicated example paragraph in the partition guard comment.
- Discard InverseWithLevel's IterMapLevel explicitly instead of
keeping an unused pair binding.
* [Test] Cover reducer accumulation inside a T.Pipelined body
Pipeline planning must classify the per-iteration update intrinsic as
compute while the shared staging copy forms the only copy stage; the
epoch spans every pipeline iteration with a single finalize. This path
previously had no coverage at all.
Warp specialization is disabled in the test: a WS-split epoch is out of
scope for reducer v2 and currently surfaces as the finalize lowering's
thread-bounds ICHECK rather than a user-facing diagnostic.
* [Cleanup] Revert the pipeline-planning reducer classification marks
The three IsReducerUpdateCall sites marked update intrinsics as compute
during pipeline statement classification. Empirically they change
nothing: with the marks compiled out, both a plain pipelined
accumulation and an adversarial mixed statement (copy-shaped shared
store plus an update in one parallel loop) produce byte-identical
kernels — a bare intrinsic never classifies as a copy on any consumer
path that matters while warp-specialized reducer epochs are rejected
outright. pipeline_planning.cc returns to its upstream state, and the
now-unused IsReducerUpdateCall predicate is dropped with it.
* [Lint] Apply pre-commit clang-format
The multi-segment recursion in tcgen05_{ld,st}_core advanced the
register pointer in dst_t/src_t elements while segment lengths count
b32 columns, so a 16-bit copy whose column count is not a power of two
(e.g. 48 columns = one x32 + one x16 segment) accessed the tail
segments' registers at half the intended offset -- a silent
miscompilation. Advance the pointer in b32 units instead.
Add asymmetric whole-store/whole-load roundtrips at (128, 96) bf16: a
symmetric roundtrip cancels the unit error, so exactly one side is
split into power-of-two (single-segment) slices, catching each
direction independently.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Aliasing buffers must induce the same storage-bit map,
bit(k) = flat_output(k) * elem_bits: layout outputs are measured in the
buffer's own elements, so a reinterpreting view needs unit conversion on
the last (stride-1) output mode, which the generic flat-index reshape
cannot express (it drops the position inside a storage slot).
TryReinterpretReshape implements the rule uniformly for every element
width, replacing the packed-sub-byte special case:
- Narrowing by f: old element q becomes new elements f*q .. f*q+f-1,
always representable as last_output * f + lane. The sub-byte
lane-append/collapse forms are deleted: the trailing lane mode was
only ever a legality certificate for the widening direction.
- Widening by f: representable only when every aligned group of f old
elements sits contiguously in one aligned slot, now proven per
reshape (plus thread-mapping invariance for fragments); layouts that
interleave storage below the new element width are a hard error
instead of silently miscompiling.
* [Enhancement] Expose cluster_mask on T.tma_copy
TMA multicast already composes with tma_copy's split-phase contract: in
LowerBulkCopy, use_multicast is orthogonal to the flag that suppresses the
wait, so annotations={"cluster_mask": ...} produces an issue-only multicast
today. Only the named parameter was missing.
Promote it to a documented keyword argument, matching the parameter that
copy_cluster already exposes under the same name, and write down the
semantics: the lowest-ranked CTA in the mask issues the multicast, other
in-mask CTAs issue nothing, a CTA outside the mask falls back to a unicast
load, and every CTA runs its own expect_tx against its local barrier.
No lowering change.
* [Testing] Drop the duplicated TMA store arrive/wait
T.copy already lowers shared -> global to tma_store + tma_store_arrive +
tma_store_wait, so the explicit pair after it emitted a second, redundant
arrive/wait.
* [BugFix][Hopper] Select the widest legal WGMMA N instead of gcd
block_N extents that are legal single WGMMA widths but not powers of two
(96, 112, 160, 176) were decomposed into n32 / n16 instructions, because
the instruction width came from gcd(warp_col_tiles, 256). wgmma.h already
instantiates every multiple of 8 up to 256, so only the width selection
was the constraint.
Select the widest legal width that tiles warp_col_tiles exactly, keeping
gcd as the fallback. Widening is restricted to multiples of 16: the
N % 16 == 8 class already miscomputes (#2593), and the integer instruction
tables only instantiate multiples of 16 (plus 8 and 24). The result is
never narrower than before, so no existing shape regresses.
The sparse generator carried the same line and now shares the helper.
* [Testing] Cover T.gemm_sp at the widened WGMMA N values
The width selector is shared with the sparse generator, so assert the emitted
wgmma_sp instruction and the sparse numeric result for block_N in {96, 112,
160} as well. Verified as a real regression: with the selector reverted the
sparse path emits n32 / n16 / n32 exactly like the dense one.
On SM70-79, normalize_atomic_type leaves bfloat16_t unchanged, but CUDA
provides no atomicAdd overload for it. Scalar, return-prev, and vectorized
bf16 atomic adds therefore failed to compile; AtomicAddRet also fell through
to unsupported cuda::atomic_ref<bfloat16_t>::fetch_add.
Implement atomicAdd(bfloat16_t*) with a 16-bit CAS loop, route AtomicAddRet
through it, and add a per-element scalar fallback for x2 that the existing
x4 path composes below SM80. Guard the overload to SM70+ so unsupported
16-bit atomicCAS does not affect older targets. Add contended numerical
coverage and an explicit sm_75 compile-only regression for the scalar,
vectorized, and return-prev helpers.
