Resolve conflicts in merge_shared_memory_allocations.cc and codegen_cuda.cc:
upstream's new shared common/storage_size.h GetBufferStorageSizeBytes helper
covers the packed-FP4 exact-bits sizing this branch introduced (and extends
it to 4-bit integer types), so the sizing conflicts take the upstream side
and the branch-local sizing helpers are dropped. The NVFP4
byte-offset -> logical-index conversion in the non-alias rewrite stays, as
upstream has no equivalent. codegen_cuda.cc's 4-bit shared-allocation
handling was equivalently refactored upstream (is_byte_packed_4bit) and
already includes the scalar float4_e2m1fn case.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
T.copy_ue4m3_scale_tile and T.ue4m3_scale_tile_source_coords were
NVFP4-specific helpers on the general T.* surface. Remove both and move the
abstraction down a level:
- The example now declares the packed scale source with its honest "tile rows"
shape ([n_mn_blocks * n_k_blocks * block_MN, block_K // 64] uint32): the
blockscaled_chunk_kmajor layout stores each (mn_block, k_block) tile
contiguously, so with this view every tile is a plain rectangular slice and
staging is a two-line T.Parallel loop with row-offset addressing only. The
host obtains the view zero-copy via
swizzle_blockscaled_chunk_kmajor_scale_words(...).reshape(-1, 4).
- tilelang.quantize.nvfp4 gains blockscaled_chunk_kmajor_tile_source_coords,
the hand-written offset formula for staging schedules that keep the 2-D
[MN, K/64] view (the maint benchmark's producer warp group uses it).
Compiler-internal staging was prototyped first (plain T.copy of the slice)
and rejected for now: the raw-pointer bulk-TMA path this lowers to on some
shapes never arrives on its mbarrier (deadlock, tracked upstream), and the
synchronous fallback stalls the producer warps (-8.5% at 2048^3). The example
comment documents this so the loop can become a plain T.copy once fixed.
Validation on RTX 5090: 123 tests pass; example outputs bitwise-identical to
the previous macro implementation on 5 shapes including M/N tail tiles
(257x384x512, 130x128x256, 128x136x256); example perf at parity
(100.1-100.4% over 2048^3/4096^3/8192^3/8192x8192x1536, interleaved A/B,
generated CUDA identical modulo a loop variable name); maint WS benchmark
kernel source byte-identical with --verify passing and perf unchanged
(641.9 / 1310.5 / 1149.5 TF at 2048^3 / 8192^3 / 8192x8192x1536).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Address maintainer review follow-ups:
- Move the four loose SM120 NVFP4 files under maint/gemm/ into a dedicated
maint/gemm/gemm_sm120/ folder to distinguish them from other gemm types,
updating the repo-root path resolution, the benchmark usage string, the
cutlass_nvf4_ref.cu source path, and the quantizer loader in the scale-layout
test.
- Guard the SM120-hardware test points with
tilelang.testing.requires_cuda_compute_version_eq(12, 0): replace the manual
capability skip in the example CLI test and tighten the block-scale kernel
tests from _ge(12, 0) to _eq(12, 0).
- Reuse the native tilelang.cuda.target.target_is_sm120 helper in
_is_explicit_non_sm120_cuda instead of hand-parsing the arch string, keeping
the explicit-arch guard so targets without an arch attr never trip the C++
ICHECK on hosts that have no SM120 GPU.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
* [BugFix] Pack 32-lane 8-bit CUDA vectors correctly
* [Test] Avoid SM100-only store in int8 broadcast regression
* [BugFix] Move int8x32 packing into CUDA templates
* [BugFix] Track GEMM accumulator writes in warp-specialization liveness collector
* [BugFix] Track scale-factor regions and region extents in gemm liveness
T.atomic_load and T.atomic_store previously accepted every order in the shared memory-order map, including operation-illegal combinations that trigger device-side assertions in libcu++.
Validate memory orders against operation-specific allow lists before lowering, report invalid values with a frontend ValueError, and update the documentation and regression coverage accordingly.
Fixes#2574
* [BugFix] Return previous value for atomic min/max/add (scalar + vectorized)
`atomic_min`/`atomic_max` (scalar) and `atomic_addx2`/`atomic_addx4`
(vectorized) typed their call with a non-void return under
`return_prev=True` but always selected the void `*_elem_op`, so codegen
emitted an empty binding (`prev = ;`) and nvcc failed.
Select the `*_ret_elem_op` variant under `return_prev`, as `atomic_add`
already does:
- min/max: the ret ops, codegen handlers, and device templates
(AtomicMaxRet/AtomicMinRet) already existed.
- addx2/addx4 (float32): register the ret ops (builtin.cc/.h) and add
codegen handlers emitting AtomicAddx2Ret/AtomicAddx4Ret (already in
the device headers) as expressions; the return is typed as the packed
2-/4-lane vector so the store matches.
- addx2 (fp16/bf16): the ret store LHS is `uint1`; codegen wraps the
half2/bf162 result with `tl::to_uint1`.
- addx4 (fp16/bf16): add AtomicAddx4Ret(half_t*)/(bfloat16_t*) device
overloads that run two AtomicAddx2Ret halves and return them packed as
uint2 (a fp16x4 is stored as uint2).
The 16-bit x4 path is per-pair atomic (two half2 atomics), not
whole-quad single-atomic -- hardware has no vector atomic wider than
`.f16x2` for 16-bit types. This matches the existing fp32-x4 scalar
fallback, and the docstring only promises "quad-width atomic addition".
Adds tests across scalar/x2/x4 and float32/fp16/bf16.
Fixes#2667
* minor fix for missing conditiions
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
`T.infinity` aborted with an internal `LOG(FATAL)` ("Cannot decide
infinity for type float8_e5m2") because `infinity_op` had no fp8 branch.
e5m2 is the only fp8 format with a representable inf, so add a branch
returning +inf, and emit it in codegen via its bit pattern
(`::bitcast(0x7c)`, mirroring the tfloat32 path) since inf/NaN cannot be
spelled as a float literal for the fp8 constructor. Other fp8 formats
(e4m3*, e3m4, e8m0, finite fnuz incl. e5m2fnuz) have no inf and keep the
fatal. Adds e5m2 to the infinity test.
* feat(autotuner): support per-config pass_configs with code deduplication
- Extract _merge_pass_configs_into_compile_args helper to deduplicate
pass_configs merge logic across grouped_compile.py and tuner.py
- Replace json.dumps uniformity check with direct equality comparison
- Extract _make_jit_compile_func factory to unify lazy/eager jit_compile
- Fix decorator-mode cache isolation: use compile() directly when
per_config_pass_configs is present to avoid JIT cache key collision
- Add example and test for per-config pass_configs tuning
* docs: add docstrings to meet coverage threshold
* refactor(autotuner): move pass_configs resolution to tuner so grouped compile handles uniform configs
* fix(autotuner): define _PASS_CONFIGS_KEY constant and fix jit_compile race condition
- Replace all magic string "__pass_configs__" with _PASS_CONFIGS_KEY constant
- Move _pass_configs_lock to cover both branches in jit_compile to prevent
concurrent reads of temporarily mutated pass_configs
* refactor(autotuner): introduce type aliases for compile-unit bookkeeping
- Add ConfigArg, UnitItem, UnitResult, CompileUnit type aliases
- Replace all long tuple type annotations with aliases
- Expand example configs to include num_stages=[1,2] and enable_grouped_compile
---------
Co-authored-by: daniu.ld <daniu.ld@alibaba-inc.com>
T.dp4a previously accepted arbitrary buffer dtypes and passed them to backend DP4A implementations, which could silently reinterpret the operand bytes and produce incorrect results.
Require both inputs to use int8 and the accumulator to use int32. Add regression coverage for valid Buffer and BufferLoad forms and for invalid input and accumulator dtypes.
Fixes#2600
* Add experimental IKET CUDA backend hooks
* Keep IKET APIs under namespace
* Use real Perfetto screenshot for IKET example
* Show IKET events in Perfetto timeline screenshot
* Add IKET validation overhead example
* Remove IKET validation overhead example
* Fix IKET lint issues
* Move IKET integration to CUDA tools
* Add tools documentation and fix workflow regressions
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
Follow-up to the package export slimming: blockscaled_chunk_kmajor_word_offset
is no longer re-exported at the tilelang.quantize package level.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
tilelang/quantize's identity is device-side conversion helpers plus
host-side data-prep utilities (quantization.py's _tir_* converters, the
lop3/mxfp intrin groups); a complete quantizer operator does not belong
in the library. Move the TileLang BF16->NVFP4 kernel factory and its
dispatcher to maint/gemm/tilelang_nvfp4_quantizer.py, next to the maint
benchmark that consumes it.
nvfp4.py now carries exactly the format contract: the _tl_* device
helpers (FP4/UE4M3 encode/decode, scale-source coordinates) any user
kernel can call, the blockscaled_chunk_kmajor packers/swizzles, and the
torch reference quantizer — 524 lines (1334 at series start). The
package re-export list shrinks to the four user-facing names; the alias
wrappers are gone.
Validated: quantize + example CLI suites (35 passed), maint benchmark
and example --verify pass end-to-end.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
nvfp4.py: remove the two remaining bring-up era quantize kernels (the
legacy no-TMA-load 16x256 schedule and the k64_per_cta=1 generic path,
reachable only through tuning knobs nothing uses) and the use_tma_load /
k64_per_cta parameters. The dispatcher now maps straight onto the single
tiled kernel factory; outputs verified bitwise-identical to the previous
code for every remaining configuration. 978 -> 710 lines (1334 before
the slimming series).
Example: drop the ones / constant-scale debug input modes and their CLI
flags; the random-FP4 + random-binary-scale path is the only
verification mode and checks exact equality.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two structural reductions, no behavior change:
- Merge the three TMA-staged quantize kernel variants (16/32/64x256)
into one _tilelang_nvfp4_blockscaled_quantize_kernel_tiled factory:
their inner loops were line-identical, differing only in rows_per_cta,
stage count and the store path (TMA store vs direct global stores),
all of which are trace-time parameters now. Outputs verified
bitwise-identical to the previous per-variant kernels for every
dispatch configuration (quant bytes and scale source).
- Move the 14 SM120 OMMA scale-selector contract oracles into the
contract test file that is their only consumer; the quantize module
stays focused on quantization and layout packing.
1334 -> 978 lines. Quantize suite: 32 passed.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Production shapes rarely keep M (batch x seq) a multiple of 128. Lift
the divisibility requirements to the validated envelope:
- scale packers zero-pad rows to full 128-row tiles (the same ceil
convention CUTLASS/cuBLAS/CuTeDSL use for the blocked SF tensor), so
the packed source always has ceil(rows/128)*128 rows;
- the example accepts an arbitrary M or an N tail: TMA loads zero-fill
out-of-bounds rows and the C store is predicated. N keeps the
16-byte-aligned-rows rule (N % 8 == 0), the same contiguous-dim
alignment CuTeDSL validates;
- simultaneous M and N tails are rejected with a clear error: the
fragment-to-global copy lowering emits the column guard bound from
the wrong dimension in that case (upstream bug, tracked separately).
Tests: tail-shape runtime correctness (257x384, 130x128, 128x136 with
exact binary-scale references), rejection paths, and packer padding
round-trip.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The example and the maint warp-specialized benchmark each carried a
private macro with the blockscaled_chunk_kmajor tile addressing spelled
out by hand. Move that knowledge into the library next to
T.mma_gemm_blockscaled:
- T.ue4m3_scale_tile_source_coords: the single source of the packed
scale tile addressing (tiles are contiguous in this layout);
words_per_kblock stays a parameter so future contracts
(block_K != 256, MX scale packings) reuse it.
- T.copy_ue4m3_scale_tile: cooperative whole-block staging macro with
trace-time dtype/shape validation, used by the example. The scale
operand stays visible in user kernels; only the layout math is hidden.
The maint kernel keeps its producer-warp-group staging schedule and
reuses the coords helper. Validated: nvf4 language + example CLI tests
(91 passed incl. a new validation test), example and WS benchmark
--verify pass.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
blockscaled_chunk_kmajor packed words are byte-identical to the
CUTLASS/CuTeDSL canonical blocked SF layout (tile_atom_to_shape_SF,
sf_vec_size=16, atom ((32,4),(16,4)):((16,4),(0,1)) tiled with order
(2,1,3)); the two views bridge zero-copy via torch view casts. Document
that on the packer and lock it with a cross-layout contract test.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Merging the SM120 block-scale mode into TensorCoreIntrinEmitter changed
mma()'s 4th/5th positional parameters to SFA_buf/SFB_buf, so the plain
gemm lowering's positional mma(A_local, B_local, C_buf, ki) call bound
ki to SFA_buf and raised 'Scale buffers require TensorCoreIntrinEmitter
block-scale mode' for every non-blockscaled MMA GEMM (CI examples:
attention_sink, deepseek_nsa/mhc/v32).
Keep the base (A, B, C, k_inner) positional signature, make all
scale-related parameters keyword-only, forward k_inner to the base
implementation, and drop the unused blockscaled mma_atom override that
shadowed the base method the same way. Add a signature-contract
regression test.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Move the persistent CUTLASS comparison harness to maint/gemm and keep the public SM120 NVFP4 example focused on a non-persistent TileLang kernel using T.mma_gemm_blockscaled and normal T.copy output.
Add a source-level guard so the public example does not grow CUTLASS build or subprocess benchmark plumbing again.
Remove the public micro_pipeline strategy knob from T.mma_gemm_blockscaled and route the packed-scale SM120 path internally through the package-pingpong lowering. Hide SM120 strategy helpers from the public TIR API and remove stale emitter experiments that were only used by the strategy sweep.
Restore the optimized persistent SM120 NVFP4 example as the benchmark path, but keep the language surface clean by using sf_layout to select the packed scale contract and a regular T.copy epilogue instead of private C-fragment store APIs.
Validation: pre-commit run --files on changed files; pytest testing/python/language/test_tilelang_language_nvf4_mma_block_scale.py testing/python/language/test_tilelang_sm120_nvfp4_example_cli.py -q => 89 passed; 512^3 correctness passed; 8192^3 128x128x256 stages=2 timing => 0.7198 ms, 1527.57 TFLOPS.
