Gfx1250/microbench - #5239
Merged
JiaoliangYu merged 81 commits intoSep 3, 2026
Merged
Gfx1250/microbench#5239
Conversation
The a8w8 rows in this file were tuned before ROCm#4151 renamed the FlyDSL kernels, and that PR retuned four CSVs but not this one. All 49 of its FlyDSL rows have failed to parse ever since: _parse_flydsl_kernel_name returns None for the old five-field name and the caller quietly falls back to the default CK kernel, so those shapes have been running untuned. Retuning is what actually fixes them; the parse-failure path is silent by design and is left for a separate change. Split-K is now in the search space (ROCm#5007), and 47 of the 134 rows pick splitK > 0. Running the op under each config on the 49 shapes -- old being the CK fallback those rows really reach today, not the kernel they name -- puts the new config ahead by 3606 -> 2702us in total, a median of 22.8%, with no shape behind by more than 0.2%. M=2 N=2624 K=6144 goes 13.9 -> 6.9us. Measuring this needs one non-obvious step. gen_instances.py compiles the tuned CSV into the lookup table that ck and cktile dispatch through, but its output is not part of the JIT build signature, so editing a tuned CSV never invalidates an existing module. Against a module built before these rows existed, all fourteen ck and cktile rows miss the table and land on rowwise_heuristic_dispatch, which returns one fixed kernel regardless of M -- the cktile rows measure 10.6-11.7us that way against the 2.6-3.9us the tuner recorded. Deleting aiter/jit/module_*.so and aiter/jit/build/module_*/ after updating a config rebuilds the table; the numbers above are from a rebuilt module. FlyDSL rows are immune because they reconstruct the kernel from kernelName at runtime. GLM-5.2 TP4 gsm8k scores 0.9704 +/- 0.0047 exact_match on both flexible-extract and strict-match, 1319/1319 answered. New shapes for both TP4 and TP8: a8w8 gains N=2688/K=6144 and N=6144/K=12288 (TP4) plus N=2048/K=2048, N=3072/K=6144 and N=3584/K=512 (TP8); BF16 gains N=160, N=256 and N=38720 at K=6144 (TP4) plus N=19360 (TP8). BF16 keeps only powers of two for M on the two widest new groups, where the intermediate sizes are not shapes the model runs. Two BF16 shapes stay untuned, M=384 and M=768 at N=256 K=6144: each carries ~10.6k FlyDSL candidates and the JIT runs out of code-region memory partway through, independently of host RAM or VRAM. They are left in the untuned CSV so a later run retries them. Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
…ROCm#5065) The 16mx4_64nx1 decode code objects returned wrong results for some combinations of context length and KV split count. Replace all three (qh16, qh64, qh128) with current builds. qh128 additionally needs the host side brought in line with the new code object: - ABI: qh128 no longer takes the legacy 288B kernarg block. Every gfx1250 MLA kernel now uses the 120B packed preload ABI, so the qh128 exception in the dispatch layer is removed. - Launch strategy: for gqa=128 the two workgroups per (batch, KV split) are now issued along x (gdx = 2) instead of along z, and z carries only the KV split id. get_meta_param's occupancy multiplier is unchanged -- the workgroup count per (batch, split) is still 2 -- so only its comment needed updating to name the new axis. Verified on gfx1250: the previously failing (context, split) combinations now match the fp32 reference at the fp8 quantization floor (cos_diff 1.4e-4..2.3e-4, no element outside a 6e-2 tolerance) for qh64 (36 configs), qh16 64nx1 (20 configs) and qh128 (29 configs), partially filled last pages included. qh8 and qh32 32nx4_3p are unaffected. Co-authored-by: Cursor <cursoragent@cursor.com>
…prefill (ROCm#4971) * feat: add gfx950 hd256 FP8 LINEAR paged-varlen asm prefill Select the PAGED_VARLEN asm kernel for gfx950 FP8 hd256 page_size=64, then fall back to CK. * fix: honor use_ext_asm and tighten paged-prefill tests Skip page64 asm off gfx950, drop the redundant page16 case, clamp empty-page seqlen_k, and use the file's FP8 threshold. * style: match FAV3 eligibility checks in batch-prefill asm Fold the -1 ladder into one compound if like fmha_fwd_v3, and restore the CK kUseGlobalLoad comment.
Every row in this config uses a power-of-two M, so an M=48 request has no tuned entry and get_CKGEMM_config pads it to the M=64 row. That row was tuned for a different width, so it is only incidentally a good fit. This adds a tuned M=48 row for eight of the nine (N,K) groups, including three that predate ROCm#5069, so the layer is covered at M=48 rather than borrowing from M=64. Tuned with --libtype all -k --shape_grouped on gfx950 (cu_num=256) in a worktree pinned to the merge commit of ROCm#5069, so the FlyDSL candidate list matches what the config is resolved against. All eight winners are FlyDSL with errRatio 0; the widest-K groups pick splitK 2 or 4, which is where most of the gain comes from. Measured against the padded-to-M=64 behaviour, three runs, per-shape median of 100 iterations after 20 warmup, one GPU: N=2048 K=2048 5.838 -> 5.080us +12.98% N=3584 K=512 4.162 -> 3.138us +24.60% N=6144 K=12288 20.505 -> 18.898us +7.84% N=4096 K=2048 6.035 -> 5.622us +6.85% N=3072 K=6144 9.643 -> 9.176us +4.84% N=7168 K=512 4.232 -> 4.067us +3.90% N=2688 K=6144 8.916 -> 8.773us +1.60% N=2624 K=6144 8.757 -> 8.784us -0.31% (within run-to-run spread) N=6144 K=4096 is deliberately left out. Its M=64 row uses a tile_m=32 kernel, and the FlyDSL candidate generator offers tile_m in {16, 48, 128, 256} for M=48 -- 32 is not among them. So the best of the 2208 candidates timed for that shape (9.943us) still loses to what padding already gives it (9.579us median), and adding the row would cost 3.07%. Leaving the shape out keeps it on the M=64 row it uses today. Its untuned entry is removed as well so a later re-run does not silently re-add the regression; it is worth revisiting if the candidate set grows a tile_m=32 variant. No existing row is modified -- the diff is eight added lines per file. Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
…Cm#5005) * perf(mha_v4): avoid copying odd-tail FP6 V inputs Signed-off-by: jcaraban <jcaraban@amd.com> * feat(mha_v4): support grouped query attention Signed-off-by: jcaraban <jcaraban@amd.com> * feat(mha_v4): add MXFP8 raw entrypoint Signed-off-by: jcaraban <jcaraban@amd.com> * docs(mha_v4): clarify grouped-query attention contract Signed-off-by: jcaraban <jcaraban@amd.com> * feat(mha_v4): add gfx942 native FP8 kernel Signed-off-by: jcaraban <jcaraban@amd.com> * fix(mha_v4): canonicalize rotated FP8 preprocessing Signed-off-by: jcaraban <jcaraban@amd.com> * refactor(bench): simplify MHA v4 quantized runners Signed-off-by: jcaraban <jcaraban@amd.com> * perf(mha_v4): deploy gfx942 XCD-swizzled kernels Signed-off-by: jcaraban <jcaraban@amd.com> * perf(mha_v4): deploy gfx942 block kernels Signed-off-by: jcaraban <jcaraban@amd.com> * fix(mha_v4): handle singleton-head rotation strides Signed-off-by: jcaraban <jcaraban@amd.com> * perf(mha_v4): deploy retimed gfx942 I8/FP8 kernels Signed-off-by: jcaraban <jcaraban@amd.com> * fix(mha_v4): deploy corrected gfx942 PV LDS waits Signed-off-by: jcaraban <jcaraban@amd.com> * fix(fmha): deploy gfx942 V staging Signed-off-by: jcaraban <jcaraban@amd.com> * fix(fmha): update gfx942 I8FP8 kernel Signed-off-by: jcaraban <jcaraban@amd.com> * feat(mha): add bf16 to mha v4 Add raw BF16/NONE dispatch and the gfx950 block kernel to the MHA v4 manifest. Generalize launcher strides to byte units, preserve the v3 aiter_bf16 benchmark, rename v4 benchmark providers to mha4_*, and cover BF16 recipe, finite output, and compiled parity. Signed-off-by: jcaraban <jcaraban@amd.com> * perf(fmha): deploy optimized gfx942 block kernels Signed-off-by: jcaraban <jcaraban@amd.com> * style(mha_v4): apply repository formatting Signed-off-by: jcaraban <jcaraban@amd.com> * test(mha_v4): isolate compile parity cases Signed-off-by: jcaraban <jcaraban@amd.com> * fix ruff warnings Signed-off-by: jcaraban <jcaraban@amd.com> * fix(mha_v4): enforce contiguous rotation layout Dense rotation kernels flatten all leading dimensions into rows, so their row stride is the last dimension width rather than stride(-2). PyTorch permits arbitrary stride metadata on singleton dimensions, which made contiguous [B, S, 1, D] inputs report a misleading head-axis stride and caused incorrect row addressing. Require contiguous dense inputs and outputs, use canonical input/output row widths, and validate output shapes, devices, auxiliary tensors, and empty inputs. Add regression coverage for singleton heads and rejected unsupported layouts. * fix(mha_v4): update deterministic BF16 kernel * Revert "fix(mha_v4): handle singleton-head rotation strides" This reverts e79b1c8 and adds rotate_activation_hd128() to mha_v4 own .cu Signed-off-by: jcaraban <jcaraban@amd.com> * Sparse MHAv4 initial commit Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * Enable sparse GQA. Fix rebase issues. Fix rotate_activations bug Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * Give MHA v4 its own hd128 rotation instead of calling into dsv4 The FP8 raw recipe rotated Q and K through module_dsv4_rotate_quant, which registers no aiter_tensor_t and so rejects the instance torch_to_aiter_pybind builds from module_aiter_core. Every mha_v4() call with an FP8 q/k format failed on that TypeError, block-sparse ones included. The MX quantizers here already run the same rotation before quantizing, so hadamard_rotate_kernel stops where they diverge and emits it in the input dtype: bitwise identical to the dsv4 kernel it replaces, and not gated on gfx950 since the FP8 recipe also runs on gfx942. A new test pins the transform against an explicit Hadamard matmul. An autouse fixture resets Dynamo per test, because the FP8 compile parity tests no longer die early and so exhausted the shared recompile limit, breaking whichever test compiled next. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * Cut the fixed cost of the sorted-sparse work table Rebuilt on every call at a cost independent of sparsity, so it came to dominate the packed call as density dropped. Two device syncs came from reading lut_count back to detect uniform counts, and thirteen ATen ops packed a few hundred elements. A stable descending sort yields the identity permutation for uniform counts without that readback, and the packing is now one kernel. On the shape measured the sparse call is ~2.6x faster at 2% density and ~1.2x at full. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * Build small sorted-sparse work tables in one kernel Order and pack the table by counting each entry's rank in LDS instead of calling ATen's sort and packing in a separate launch. The key packs the LUT length with the slot index, so ranks are distinct and stable by construction, and the low half is already the permutation the packing needs. Build time at 512 entries drops from ~23us to 9us. The quadratic rank count loses to ATen past ~1024 entries, so larger tables keep the sort path. Also expose the builder and test its ordering. A wrong order only unbalances the waves rather than changing the result, so no attention test can see it. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * Rank work table entries across a wave instead of a thread Each entry's rank was counted by a single thread walking every key, which costs O(n) per thread and lost to ATen's sort above about 1024 entries. Split the count across a wave and reduce it, so per-lane work is n/64 and the build stays near 6us from 512 entries to 4096. That moves the fused cutoff to 8192, which is where a workgroup's 64KB of LDS runs out. Wan 720p self-attention at 5 heads needs 1480 entries and so was on the fallback branch at 25.7us; it now builds in 5.9us, taking the whole call from 176us to 157us at 1.6% density. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * fix(mha_v4): restore BF16 dense dispatch and launcher byte strides Re-hook mha_v4() through mha_v4_packed for BF16/NONE, reject sparse BF16 explicitly, and pass byte strides (skipping descale setup) in populate_dense_kernarg. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * style: run black, ruff, and clang-format on block-sparse MHA v4 changes Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * docs(mha_v4): trim sparse section implementation detail Drop kernarg offsets, bit-packing formulas, and duplicate sparse GQA text from mha_v4.md; keep API contracts and move sparse GQA notes into Sparse Contract. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> * Add gfx942 sorted-sparse MHA v4 kernels The gfx942 FP8/FP8 and INT8/FP8 sparse rows use a 256x64 tile rather than gfx950's 256x128, so sparse geometry is no longer arch-invariant: mask shapes go through mha_v4_kv_tile(), and the key-length check reads cfg.ts_kv instead of a literal 128. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * fix(mha_v4): guard sparse launches by device, validate LUT contents The sparse launcher installed HipDeviceGuard after build_sorted_work_table, whose raw HIP kernels take the current device rather than Q's, so a launch with Q on a non-current GPU faulted; mha_v4_sparse_work_table had no guard and silently returned zeros. Move the guard above every device query and launch, and add one to the work-table op. Also reject non-bool and wrong-device block_mask, bound kv_block_indices against the row count, and add opt-in AITER_MHA_V4_VALIDATE_LUT=1 for device-side checks. Empty LUT rows fault in the ASM rather than acting as no-ops, so document them as invalid. Add tests proving sparse selection follows the LUT per tile, per head, and across query tiles. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * fix(mha_v4): make an empty sparse LUT row write zeros lut_count == 0 faulted the sorted-sparse ASM, so the launcher declared empty rows illegal. Rebuild the ten sparse code objects with the prologue reads clamped and the row's KV traversal skipped, then follow through on the host: drop the kLutEmptyRow rejection, and relax the unconditional kv_block_indices bound, which was derived from the one-block-per-row assumption and would now reject a valid tightly-packed LUT. The buffer must still be non-empty, since the kernels dereference the row base even for a row that selects nothing. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * perf(mha_v4): rebuild the gfx942 i8fp8 sparse object without the hot-path clamp Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * test(mha_v4): cover a partial query tile with an empty sparse row Every sparse case used a whole number of 256-row query tiles, so the tail masking the empty-row no-op is built on was never exercised alongside a short tile. Add one case at 64/128/200 trailing rows that checks the short tile still reads the KV blocks its row names and that an all-False row on it returns zero. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * fix(bench_sage): pass the sparse LUT for the MX MHA v4 recipes mha4_mxfp4/f4f4/mxfp6/f6f4 called mha_v4_packed directly instead of the launch_mha_v4_packed wrapper that injects the LUT kwargs, so --block-sparsity was silently ignored and every density measured dense. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * refactor(mha_v4): rebuild the gfx950 f4f4 sparse object with a prologue-only clamp Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> * mha_v4: take the sparse KV tile from the manifest mha_v4_kv_tile() restated ts_kv as per-arch constants while the launcher read it from the manifest row it dispatches on. Read the CSV instead (mode=1 rows), behind torch_compile_guard since Dynamo traces a cached body and open() broke fullgraph on the block_mask path. Adds the compile test, and moves the work-table build measurements into mha_v4.md. Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> --------- Signed-off-by: jcaraban <jcaraban@amd.com> Signed-off-by: Niko Säkkinen <niko.sakkinen@amd.com> Co-authored-by: jcaraban <jcaraban@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
* [DCP] Enable fused indexer QK preparation * format * fix uncondition clamp * modify case
* deploy mxfp6 gemms * fix mxfp6 accuracy * fix mxfp6 source formatting Remove trailing whitespace so the clean branch passes git diff checks. Co-authored-by: Cursor <cursoragent@cursor.com> * Add MXFP6 GEMM tuning and shape-based dispatch * refactor * fix * ruff * replaces per-element log2/exp2 encoding with mathematically equivalent piecewise E2M3 encoding. * co-pilot comments fix * improved hip quantization * fix MXFP6 backend and buffer validation * Apply suggestions from code review Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * fix broken copilot suggestions * fix A6W6 ASM default kernel selection * cover all A6W6 kernels and padding paths in CI --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…Cm#4659) * [triton] Add two fused ops for diffusion transformer blocks A DiT block spends its non-GEMM, non-attention time in two patterns that torch runs as long chains of elementwise ops. Both are memory bound, and both are dominated by temporaries the maths does not need. fused_rmsnorm_indexed_adaln out[m] = rmsnorm(x[m], weight) * (1 + scale[idx[m]]) + shift[idx[m]] Adaptive layernorm: every token indexes a small table of modulation vectors, one row per (modality, timestep). Unfused, the normalised activation is written and immediately re-read, and both table gathers are materialised at [M, N] -- 680 MB each at a 63k-token request. One program owns a block of rows and walks each row in column tiles, once to accumulate the sum of squares and once to normalise and modulate, so x is read once and out written once. Two details that matter for this workload. Rows are tiled rather than padded to the next power of two, because a 5376-wide row would mask off a third of every access at 8192. And a block of consecutive tokens usually shares one modulation index -- packed sequences are laid out in runs of one modality -- so the kernel checks for that and collapses the [BLOCK_M, BLOCK_N] gather to a single [BLOCK_N] load broadcast in registers. fused_qk_norm_rope_cached q[t, h] = rope(rmsnorm(q[t, h], q_weight), cos_sin_cache[t]) (and k) Per-head RMSNorm followed by partial NeoX RoPE, on q and k, in place. The existing rope ops do not cover this case: they assume the rotated subspace is the whole head or half of it, and diffusion transformers rotate fractions in between (96 of 128 for MiniMax-H3), while the cache-write variants want a paged KV cache that a diffusion model does not have. One program owns a token. A token's heads are contiguous, so the [H, D] tile is one coalesced run and the token's cos/sin row is read once for all heads rather than being broadcast into a [T, H, D] temporary. Only each token's [H, D] block must be contiguous, so q and k can be strided views into a packed qkv projection and are rotated in place, never materialised. Measured on MI355X, bf16, at MiniMax-H3's shapes: rmsnorm + indexed adaln, 63232 x 5376 1.708 ms -> 0.329 ms 5.2x qk norm + rope, 63232 x 56 x 128 11.690 ms -> 1.262 ms 9.3x Both hold their speedup across the token counts one rank sees at Ulysses 1/2/4/8. Accuracy: both keep the row in fp32 across the whole fusion, so they are nearer the fp32 result than the op chain they replace, which rounds to bf16 at each step. Against that chain on a real 50-layer model, one forward agrees to cosine 1.0000000 with max relative error 1.3e-4. 58 tests: every table row exercised individually (a kernel that broadcast row 0 would pass a uniform-index test), the uniform and scattered index paths checked against each other, q and k given different norm weights, the unrotated tail checked for passthrough, and the strided-qkv-view case checked to leave v untouched. * Address review comments on the diffusion adaLN / RoPE fusions Test fixes: - test_uniform_and_scattered_indices_agree asserted nothing. Both index tensors were torch.full((M,), 2), so `fast` and `slow` were the same call on the same input and assert_close(atol=0) could not fail. That left the kernel's `uniform = tl.min(idx) == tl.max(idx)` branch -- which broadcasts one modulation row instead of gathering [BLOCK_M, BLOCK_N] -- with no coverage at all. The intent was also unreachable as written: with a single index value no arrangement is ever non-uniform. Two table entries are now made identical, so the same modulation is reachable both uniformly (broadcast branch) and alternating (gather branch) and the two must agree bit for bit. Verified by breaking the uniform branch on purpose (broadcast table row 0 rather than the block's index): the old assertion still passed, the new one fails. - Every call whose output is asserted on now states `eps=1e-5` rather than leaning on the wrapper default, matching the `reference` calls beside them. The default is that same value, so nothing was computing the wrong thing, but the tests should not depend on it staying put. The two `pytest.raises` calls keep the default, where eps plays no part. Kernels and wrappers: - 1.0 / tl.sqrt -> tl.rsqrt in both kernels. Checked rather than assumed: all 58 tests pass unchanged, including the fp32 cases at 2e-6 / 2e-5. - Lazy %-style logging instead of eagerly built f-strings. This needed AiterTritonLogger to forward *args -- its methods took (self, msg) only, which is why the f-string pattern is everywhere in the Triton kernels. The change is additive, so existing single-argument callers are unaffected. - The one assert in the adaLN wrapper without an error message now has one. - Dropped the `if M == 0` / `if T == 0` guards. They were speculative; no framework hands these ops empty tensors. - The RoPE wrapper docstring showed only the q equation; k was missing.
…estigial signal/buffer block, add explicit raw-pool override (ROCm#4924) * fix(dist): init_dist_env no longer breaks raw IPC input pools Under PYTORCH_HIP_ALLOC_CONF=expandable_segments:True -- the very configuration the raw_cached input pool exists for (ROCm#4174) -- init failed twice over in init_dist_env's signal/buffer block (ROCm#4921): * register_input_buffer(signal) exports the signal tensor's pointer via hipIpcGetMemHandle, but the torch.zeros signal is VMM-backed under expandable segments and the export dies at custom_all_reduce.cu:417 with 'invalid argument'; * ca_comm.buffer = ca_comm._pool["input"].tensor raises by design, because the raw_cached pool is a plain hipMalloc region with no backing torch.Tensor. The block is removed rather than repaired, because all of it was vestigial: * ca_comm.signal / ca_comm.buffer are never read anywhere in the tree; * C++ register_input_buffer only inserts a pointer-translation entry keyed by the registered tensor's own address, which is consulted only when an allreduce is invoked with that exact tensor as input -- something that never happens for the signal tensor (open_ipc_handle's handle cache is filled on demand, so no pre-warming is lost either); * gfx1250 has skipped the entire block since its VMM bring-up (the vmm_exchange rendezvous deadlocks) and works without it. CustomAllreduce.__init__ already builds its own meta/input pools and forces the copy-in path under expandable segments, so nothing here was load-bearing. get_tp_group stays imported: this module is a re-export surface (downstream engines import set_custom_all_reduce through it). Adds op_tests/multigpu_tests/test_init_dist_env.py: brings up init_dist_env per rank under both allocator modes (default torch pool, and expandable_segments -> raw_cached) and checks one allreduce. The existing test_custom_allreduce.py performs its own init and never executes init_dist_env, which is how the regression shipped. Fixes ROCm#4921. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * feat(dist): AITER_CUSTOM_AR_RAW_INPUT_POOL forces the raw IPC input pool The raw (plain-hipMalloc) input pool previously had exactly one trigger: PyTorch expandable segments. But expandable segments break custom allreduce later anyway -- every capture-time output is a fresh torch.empty_like whose VMM pointer get_output_buffer_RD records for post-capture IPC export, which then fails in get_graph_buffer_ipc_meta -- so the raw pool's one trigger leads to an unusable configuration (ROCm#4921, third failure mode; ROCm#4621's copy-in guard covers inputs only). The override gives the raw pool a trigger that works: co-resident engines on one node, where a second engine's torch.empty input pool can fail hipIpcGetMemHandle outright. Under the default allocator everything else (meta pool, capture-time outputs, graph flush) stays exportable, so only the input pool needs to move to hipMalloc. Extends test_init_dist_env.py with a raw_override mode that asserts the flag actually selects the raw pool and allreduce stays correct. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * feat(dist): log the input-pool allocation mode at init A silently-inert pool trigger is indistinguishable from a working one by behaviour alone -- the engine serves fine single-engine either way, and the failure modes this pool exists to avoid (ROCm#4921) only appear in specific modes under specific co-residency. One INFO line per rank makes every run self-document which pool it actually got, so a mislabeled measurement is catchable from the log rather than by re-deriving the allocator state. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * fix(dist): honor capture registration setting in fused AR --------- Co-authored-by: ThomasNing <thomas.ning@amd.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
… for T512 (ROCm#5070) * perf: TDM prefill bandwidth opt — K=5 occupancy + position prefetch Two changes to improve TDM prefill kernel bandwidth at small T: 1. Reduce LDS buffer count from K=6 (192KB) to K=5 (160KB) when num_rows <= 65536. On gfx1250 with 320KB LDS per CU this allows 2 WGs/CU instead of 1, doubling occupancy. 2. Prefetch position buffer_load before the hot loop: issue the first group's position load before TDM prologue, and each subsequent group's position load after the prior group's last tile compute. This overlaps the position→cos/sin serial dependency chain with TDM tile transfers and compute, reducing loadcnt stalls by ~41%. ATT trace confirms total stall cycles drop 36% (72K → 46K), with loadcnt (HBM) stalls down 41% and dscnt (LDS) stalls down 70%. Measured kernel times (gfx1250, H=128 D=512 RD=64 BF16): T=512: 16.5us → 11.9us (+39%) 8.2 → 11.4 TB/s T=16384: 302us → 274us (+10%) 14.3 → 15.8 TB/s Co-Authored-By: Claude <noreply@anthropic.com> * perf: drop TDM prefill rotation to K=4 at small T At num_rows <= 65536 (T=512, H=128) CT=8 yields gx_q=256 workgroups for 256 CUs -- exactly one WG per CU. LDS is therefore never the limiter at this shape, which invalidates the reasoning behind the previous K=5 choice (it was picked to keep the arena at 160 KB so two WGs would fit, but a second WG never exists here). With K free to pick on latency alone, K=4 measures faster. T=512 H=128 D=512 RD=64 BF16, gfx1250, three runs each: K=5: 15.728 15.647 15.683 -> 15.69 us (8636 GB/s) K=4: 15.357 15.360 15.220 -> 15.31 us (8874 GB/s) Non-overlapping ranges, ~2.4% faster. T=16384 is unaffected (it takes the num_rows > 131072 branch at K=6): 302.6 / 307.5 us, unchanged. The mechanism behind the shallower rotation winning is not understood -- it is not LDS or occupancy driven -- so the docstring records the measurement and warns against extrapolating to other shapes. Also measured and rejected on this shape: - CT=4 to reach 2 WG/CU: 14.97 vs 14.86 us, no gain. Doubling the wave count doubles the per-wave cold-start cost, cancelling the extra latency hiding. - Issuing the position load before the TDM prologue: 15.63 vs 15.68 us, within noise. The K descriptor setups are far too few instructions to cover a ~1700-cycle DRAM miss. - TDM store (LDS -> global) in place of buffer_store, tried with a reused input buffer, one dedicated output buffer, and two rotating output buffers: 16.04 vs 15.27 us at matched K=4, ~5% slower. The LDS round trip (ds_write plus tensorcnt sync) costs more than the s_wait_xcnt it removes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: erratum for 53d6700 — its perf claims and attribution were wrong 53d6700 ("perf: TDM prefill bandwidth opt — K=5 occupancy + position prefetch") is already published, so its message is left in place and corrected here instead. Four claims in it are wrong. The code it shipped is fine and is kept; only the reasoning and the numbers were bad. 1. "Prefetch position buffer_load before the hot loop ... reducing loadcnt stalls by ~41%" The prefetch is a no-op. issue_pos() is followed immediately by the trunci that consumes it, in the same statement, so there is no distance for the load to cover its miss: pending_pos[0] = issue_pos(tok_of(tile_base + i + 1)) cs_cache[0], cs_cache[1] = _cs_from_pos( fx.Int32(pending_pos[0].trunci(i32))) An ATT capture of the shipped code shows group 1's position still stalling 5249 cycles despite being "prefetched". The loadcnt reduction came entirely from K=6 -> K=5. 2. "K=5 (160KB) allows 2 WGs/CU instead of 1, doubling occupancy" At num_rows=65536 (T=512, H=128), CT=8 gives gx_q=256 workgroups for 256 CUs, so a second WG per CU never exists and LDS was never the limiter. 034220f already replaced this reasoning in the _tdm_tiles_per_wg docstring. 3. "T=512: 16.5us -> 11.9us (+39%)" The 11.9us came from a hand-rolled L2-warm timing loop and is not comparable to the 16.5us op_test figure it was subtracted from. Measured on one path (op_tests/test_flydsl_qk_norm_rope_quant.py), idle GPU, five runs each: K=6 16.333 16.434 16.386 16.463 16.441 -> 16.41 us K=5 15.674 15.688 15.618 15.668 15.639 -> 15.66 us (+4.6%) K=4 15.332 15.227 15.281 15.239 15.422 -> 15.30 us (+2.3%) So 53d6700 was worth +4.6%, not +39%, and the two commits together are worth +6.8% (16.41 -> 15.30 us). 4. "T=16384: 302us -> 274us (+10%)" 53d6700 does not touch that path. For num_rows > 131072, _tdm_tiles_per_wg returns (TILES_PER_WG, NUM_BUFFERS) = (40, 6), identical to the pre-commit default of CT=40 with the builder's num_buffers=NUM_BUFFERS. T=16384 measures ~305 us both before and after; the reported gain is spurious. Root cause of 3 and 4: numbers from two different timing harnesses were compared against each other. Only same-harness, same-session, repeated measurements are used above. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf+fix: TDM prefill — 8-wave workgroups, and tighten the drain-phase wait Two changes, found together while investigating why smaller workgroups appeared to help. 1. Correctness: the drain phase under-waited on its TDM loads. Tile i consumes TDM load #i; loads are issued in tile order, K in the prologue then one per iteration while i + K < CT. In steady state K+i are outstanding, so tensor_wait(K-1) leaves exactly #0..#i retired -- correct. Once the issues stop, the issued count freezes at CT and K-1 is too loose: #i is only guaranteed retired with at most CT-1-i left, which has to reach 0 on the last tile. The wait is now min(K-1, CT-1-i); both operands are compile-time constants in the unrolled loop, so this costs nothing. This was latent, not new. With ROWS_PER_TILE=32 the per-tile compute happened to outlast the load, so the shipped kernel got away with it. Shrinking the tiles exposed it: at RT=8/CT=16 the output was wrong in exactly the last three tiles of every workgroup (tile%CT histogram [0]*13 + [53,53,49]), with all 512 columns of those rows wrong -- i.e. the LDS input itself, not the RoPE tail. err_q 0.027 -> 5.96e-08 with the fix, same config, same build. Measured cost at the shipped shape (T=512 H=128, three runs each): without: 15.389 15.254 15.281 with: 15.308 15.393 15.304 2. Perf: ROWS_PER_TILE 32 -> 8, and CT 8 -> 16 for num_rows <= 65536. gx_q = num_rows / (ROWS_PER_TILE * CT) has to stay at or above the 256 CUs. At the low end of the TDM range it did not: num_rows=32768 (T=256, H=128) gave gx_q=128, so half the CUs sat idle. RT=8 restores full coverage there and doubles it at num_rows=65536. T=512 H=128, five runs each: RT=32: 15.295 15.233 15.270 15.344 15.305 -> 15.29 us RT=8: 15.090 14.937 14.929 15.009 14.936 -> 14.98 us (-2.0%) T=256 H=128, four runs each: RT=32: 12.610 12.587 12.581 12.645 -> 12.61 us RT=8: 10.904 10.499 10.587 11.065 -> 10.76 us (-14.6%) Across the TDM path (H=128 unless noted): T=256 -16.3% T=512 -2.5% T=1024 -6.5% T=16384 -2.2% T=16384 H=16 -0.9% Shapes below TDM_MIN_ROWS=32768 take the r32_w32 path and are untouched by ROWS_PER_TILE; the +-1-3% seen on those in a sweep is run-to-run noise. At RT=8, GROUP = H/RT = 16 and TILES_PER_WG=40 is not a multiple of it, so cos/sin hoisting turns off for the largest shapes. That is not a regression -- T=16384 still improves -2.2% -- consistent with the separate finding that the position->cos/sin chain is worth ~2.6% of wall clock despite being 36% of stall cycles. Validated on 20 (T, H, q_weight) combinations plus the SWA direct/paged and decode paths: all err_q/err_kv <= 1.3e-06, 16/16 checkAllclose pass. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(flydsl): optimize qk norm rope decode * perf: fuse FP8 quant into TDM prefill and tune gfx1250 occupancy Keep 2 WGs/WGP on the T=512 path, use 16-row tiles only for short prefill, and stream FP8 (grouped/e8m0) through the TDM kernel so Q write traffic drops without falling back to the slower direct path. Co-authored-by: Cursor <cursoragent@cursor.com> * perf(flydsl): TDM reads KV strided, and drops to K=2 on deep grids Two independent changes to the gfx1250 TDM path. 1. Read KV with a row stride. The TDM kernel indexed KV as `tok * D`, so the wrapper had to force kv.contiguous(). The V4 call site slices KV out of a wider qkv_a tensor, so that fired a full elementwise copy kernel on every invocation. get_trace_perf sums all device kernels, so the copy landed inside the number the op-test reports: 3.52us on top of a 12.20us kernel at T=512, 22% of the reported total, for nothing -- fused-kernel time is identical whether KV arrives strided or contiguous. Thread kv_in_row_stride through, matching what the wave32 and wave64 paths already do. 2. TDM buffer depth K=6 -> K=2 from num_rows >= 131072. K sets the length of the load-only prologue. Once the grid is deep enough that one workgroup's prologue overlaps another's steady state, the shallow K=2 wins; below that a workgroup must cover its own load latency and the deeper prologue pays for itself. Measured on top of 46ee44b, public API, rotate=4, interleaved medians: num_rows 32768 (T=256) K=2 +31.9% 131072 (T=1024) K=2 -3.0% 49152 (T=384) K=2 +9.8% 262144 (T=2048) K=2 -5.1% 65536 (T=512) K=2 +4.1% 1048576 (T=8192) K=2 -4.3% 2097152 (T=16384) K=2 -2.9% Note the crossover sits above T=512: with 46ee44b's occupancy tuning in the base, K=6 is now correct for the decode shape, so the T=512 gain here comes from (1) alone. T=512 on this gfx1250, same harness and rotation as the pre-change baseline: qw off 15.72us -> 11.07us qw on 16.48us -> 12.16us The runperf script's own sweep reports 10.74us / 10.83us for the same two rows; it launches through a tighter loop, so treat the pair above as the comparable figure. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * test: runperf script sets PYTHONPATH and prints a combined summary table `python op_tests/foo.py` puts op_tests/ on sys.path rather than the repo root, so `import aiter` failed unless the shell already exported PYTHONPATH. Set it from the script's own directory. Also tee both sweeps to a log and replay every markdown table at the end under its original heading, so the T=16384 and T=512 runs can be compared without scrolling back through two sweeps of output. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * test(flydsl): make the qk_norm_rope %peak column arch-aware _PEAK_BW_GBPS was a single 22000.0 labelled "MI355X HBM3e peak", but 22 TB/s is the gfx1250 figure -- MI355X (gfx950) is 8 TB/s and MI300X (gfx942) is 5.3. The column was therefore only meaningful on gfx1250, and silently wrong on the other two archs the file already lists in SUPPORTED_GFX. Look the peak up per arch instead. Unknown archs report None rather than a fabricated percentage. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * style: apply black to qk_norm_rope_quant.py CI runs black[colorama]==26.5.1 and this file was the only one in the branch it wanted to reformat. Formatting only -- verified the AST is identical before and after. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * chore: drop the local perf harness from the repo root runperf-qknormrope-bs16-t16384.sh is a personal benchmark driver for one shape on one machine, not something the repo should carry at its root. It is kept locally alongside the other measurement tooling. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(flydsl): let the TDM path take the fused SWA write `use_tdm` excluded kv_write and paged outright, so any caller that passes swa_kv fell back to the wave32 kernel. That is what the model does, so the decode shapes were running qk_norm_rope_H128_D512_RD64_kvw_r32_w32_flydsl and none of the TDM tuning reached them. The op-test did not show this: its headline rows pass no swa_kv, and its SWA sweep is pinned to T=8..96 by the paged fixture's capacity, so it never reaches a TDM-eligible size. Port the scatter into the TDM kernel's KV path. The gates are copied from the wave32 sibling unchanged -- bid<0, pos<0, paged blk past the table, table entry -1, resolved row past the pool -- and the row index is widened to 64 bits before the byte multiply, as there too. gfx1250, public API, rotate=4, interleaved medians: wave32 TDM gain direct T=512 17.65us 11.61us -34.2% T=1024 32.61us 23.59us -27.7% paged T=512 17.46us 11.39us -34.8% T=1024 32.07us 23.72us -26.0% Verified against the wave32 path at T=512/1024 in both modes: the pool is byte-identical to kv_out at every resolved row, rows nobody targets stay zero, and guard rows either side of the pool are untouched. Each skip gate is covered by its own case. Not covered: pos<0. It is not a legal input -- the main path indexes cos/sin with the raw position long before the scatter -- and the wave32 path faults on it identically, so this is not a new exposure. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(flydsl): address q_out past 4 GiB in the TDM path A buffer descriptor's num_records is 32-bit, so one descriptor reaches 4 GiB. q_out crosses that at T*H*D*2 >= 4 GiB -- T>=32768 at H=128, D=512 -- and every row past the limit was dropped or wrapped. The failure started exactly on the boundary: at T=32768 the first bad row was 4194303, whose last byte sits at 0xFFFFFFFF, one past num_records, so precisely 4 elements were lost. Beyond 4 GiB it degraded fast -- 0.05% of q_out wrong at T=32776, 38.9% at T=40960, NaN in both. Bias the descriptor base per workgroup instead, the same trick the SWA scatter in this file already uses. A workgroup owns CT*RT rows, so the 32-bit offset then spans 128 KB rather than the whole tensor. The bias is computed once per workgroup, outside the tile loop. This predates the TDM work: the wave32 path fails identically at these sizes, it is simply unreachable there now that TDM covers num_rows >= 32768. The op-test's default sweep includes T=65540 and had been failing on it. T=65540, H=128: err 1.4e-07 (was garbage), 12524 GB/s T=40960: 0 bad elements (was 1.04e9) No measurable cost -- T=512 11.40us, T=2048 40.66us, T=16384 330.25us, all within run-to-run spread of the previous commit. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * style: collapse the SWA store guard (ruff SIM102) The `do_swa is not None` guard was redundant -- None is already falsy, so the two ifs fold into the one the wave32 path next door already uses: if const_expr(kv_write) and do_swa: Short-circuiting still keeps the const_expr and the runtime predicate apart: kv_write=False never evaluates do_swa, emit_q passes None so no store is traced, and emit_kv passes the predicate so scf.if is emitted as before. Re-ran the SWA scatter checks (10/10) and the op-test (465 passed) to confirm codegen did not shift. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(flydsl): halve the TDM workgroup when the grid is too thin RT sets the workgroup size and therefore how many of them the grid holds: gx_q = num_rows / (RT * CT). Below 2 workgroups per CU there is no neighbour whose steady state can cover a workgroup's own load latency, and RT=8 sits under that line for num_rows < 65536 -- at num_rows=32768 it yields exactly one workgroup per CU. Halving RT there doubles the grid and pays for itself. It stops paying at 65536 (exactly 2 WGs/CU, a wash) and turns negative past it, where the smaller workgroup costs more than the extra parallelism returns (+2.4% at num_rows=262144), so RT=8 holds from 65536 up. This is the same question K already answers, one level up: can a workgroup's latency be hidden by a neighbour, or must it cover its own? gfx1250, public API, rotate=4, interleaved medians, before -> after: decode (fused SWA) prefill (no SWA) T=256 6.82 -> 6.65 T=256 6.85 -> 6.66 T=384 9.55 -> 8.53 T=384 9.30 -> 8.76 T=512 11.47 -> 11.45 T=512 11.16 -> 11.16 T=2048 41.34 -> 41.27 T=2048 41.23 -> 41.19 T=16384 333.53 -> 334.22 T=16384 331.43 -> 334.31 T>=512 is untouched by construction -- the geometry it selects is unchanged, so those rows are noise. Prefill only reaches RT=4 on prompts shorter than 512 tokens, where it is also a win. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * style(flydsl): trim the comments this branch added The tuning rationale had grown into paragraphs sitting on top of two-line functions. The measurements behind each threshold are in the commits that introduced them, so the source only needs to say what the knob does. Also folds emit_kv's inlined position load back into a load_pos() helper that load_cs() now shares. Net -26 lines. No behaviour change: op-test 465 passed, SWA scatter checks 10/10, T=512 decode 11.43us. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Satya Nikhil Kodukula <nikhil.kodukula@gmail.com> Co-authored-by: Cursor <cursoragent@cursor.com>
* Add lumen triton kernels and custom ops (clean cherry-pick) Cherry-pick of 4e19b8e (ZhangDanyang-AMD) onto upstream/main. Only new files preserved; upstream existing code left untouched. Adds: triton quant kernels, FP8/MXFP8 attention, MoE GEMM variants, cross_entropy, fused_norm_quant_gemm, AOT precompiled kernels, moe_sorting test cases. Registers cross_entropy and mxfp8_attention in triton __init__.py. * add large-M/small-N RMSNorm backward specialization * add gfx942 (MI308X, 80CU) blockscale bpreshuffle GEMM configs * add requant_fp8_row_to_col, chunked cross-entropy, add gfx942 per-shape GEMM configs compile_ops type-check fix omitted — upstream already has _is_tensor_like fix. * add gfx942 preshuffle GEMM configs for llama2-7b/13b/70b and qwen3-8b * add MoE weight gradient Triton kernel (moe_wgrad) Adds a fused Triton kernel for MoE weight gradients that operates directly on sorted_token_ids/expert_ids from moe_align_block_size, eliminating the need for sort+pad+bmm and CPU-GPU sync in backward. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> * add is_cdna4() arch probe for gfx950 family Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> * add DSV4 sparse MLA training and indexer ops for DeepSeek-V4-Flash - Sparse MLA: fused Triton fwd/bwd kernels with CSR-based dKV gather (no atomics) - Indexer: BLAS-based scoring via torch.einsum (hipBLASLt) + PyTorch autograd - Correctness tests: 84 sparse MLA tests + 48 indexer tests, all passing Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> * add Triton MHC forward and backward support Co-authored-by: Cursor <cursoragent@cursor.com> * integrate SonicMoE pure-Triton grouped GEMM MoE with full autograd Port SonicMoE's pure-Triton MoE implementation from sonic-moe into aiter-lumen. Provides trainable MoE layer with fused router + grouped GEMM + activation, supporting forward and backward passes for all 7 activation types. New files: - _triton_kernels/moe/sonicmoe/: 9 kernel modules (grouped GEMM, activations, routing metadata, reduction, forward/backward autograd functions) - aiter/ops/triton/sonicmoe.py: public API wrapper - configs/moe/gfx942-MOE-SONICMOE-BF16.json: autotune configs for MI308X - op_tests/test_sonicmoe.py: correctness + benchmark tests Correctness verified on MI308X (T=64, H=128, I=64, E=4, K=2, BF16): | Activation | output rel err | dx rel err | dw1 rel err | dw2 rel err | Status | |------------|---------------|------------|-------------|-------------|--------| | swiglu | 0.0097 | 0.0132 | 0.0138 | 0.0104 | PASS | | geglu | 0.0014 | 0.0089 | 0.0100 | 0.0000 | PASS | | reglu | 0.0014 | 0.0103 | 0.0098 | 0.0000 | PASS | | gelu | 0.0014 | 0.0134 | 0.0140 | 0.0000 | PASS | | relu | 0.0014 | 0.0155 | 0.0168 | 0.0000 | PASS | | silu | 0.0014 | 0.0146 | 0.0150 | 0.0000 | PASS | | relu_sq | 0.0014 | 0.0104 | 0.0117 | 0.0000 | PASS | All relative errors < 2%, well within BF16 tolerance. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> * fix tests: call existing topk_softmax and drop redundant RMSNorm 65536x128 The cherry-picked pytest imported a non-existent softmax_topk API; retarget it at ASM topk_softmax. 65536x128 duplicated 16384/364800 coverage of the large-M/small-N bwd path. Co-authored-by: Cursor <cursoragent@cursor.com> * fix style: format PR Python with Black and satisfy Ruff 0.16 Unblocks Checks so check-signal can let HIP/Triton CI run. Also add missing torch/_get_activation_from_str imports in gemm_a16w16_agnostic. Co-authored-by: Cursor <cursoragent@cursor.com> * drop files already removed on origin/main instead of resurrecting them Cherry-picks had re-added pre-ctypes pybind/headers, AOT hsaco, and a redundant bpreshuffle tuner. Keep gfx942 rows in the existing CSV. Co-authored-by: Cursor <cursoragent@cursor.com> * move gfx942 GEMM tunes into nested config layout so they actually load Place llama2-7b/13b/70b, llama3-8b qkv, and qwen3-8b N/K tables next to each family's DEFAULT.json. Legacy configs/gemm/ paths are ignored once the nested default exists. Co-authored-by: Cursor <cursoragent@cursor.com> * load gfx942 SonicMoE JSON at launch instead of autotuning those kernels Pick N/K/E and H buckets from {arch}-MOE-SONICMOE-BF16.json so production shapes skip the autotune search; fall back to the old autotune lists when the file is missing. Co-authored-by: Cursor <cursoragent@cursor.com> * fix ruff C408 in SonicMoE launch kwargs Rewrite dict() calls as literals so Checks reviewdog stops failing the PR. Co-authored-by: Cursor <cursoragent@cursor.com> * format PR mxfp8/moe GEMM modules for Black 26 Remove extra blank lines after module docstrings so psf/black@stable in Checks passes on CI. Co-authored-by: Cursor <cursoragent@cursor.com> * drop gfx942 CK GEMM row that duplicates DSV4 opus tune Merge keys are gfx/cu_num/M/N/K, so ck vs opus for 2048x4096x1024 on 80 CU fails wheel prebuild. Keep the faster opus entry from the DSV4 table. Co-authored-by: Cursor <cursoragent@cursor.com> * load SonicMoE JSON from nested gfx942/triton/moe layout Co-authored-by: Cursor <cursoragent@cursor.com> * format sonicmoe_config_utils for Black 26 Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: ZhangDanyang-AMD <danyzhan@amd.com> Co-authored-by: leiwu0812 <leiwu0812@users.noreply.github.com> Co-authored-by: Claude Opus 4 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com>
…0) (ROCm#5109) Mythos scan finding SEC-00830 (ROCM-26711) flags GitHub Actions event context interpolated straight into `run:` shell blocks, where the value is pasted into the script text before the shell parses it. aiter-test.yaml already uses the `env:`-indirection pattern in most steps (15 `env:` blocks; `${GITHUB_EVENT_NAME}` at lines 50/537). This brings the five remaining spots in line: - 3x `if [ "${{ github.event_name }}" = "schedule" ]` -> `${GITHUB_EVENT_NAME}` (GitHub's built-in, same as lines 50/537) - 2x `BASE_REF="${{ github.event.pull_request.base.ref || github.ref_name }}"` -> hoisted into a step-level `env:` block After this change no `${{ github.event* }}` remains inside any `run:` block. Note this is hardening, not a fix for an exploitable bug. The scanner's stated attack surface (`github.event.pull_request.title`) does not appear in any `run:` block. Of the five occurrences, three were `github.event_name` (an enumerated value) and two were `base.ref` — the PR's *target* branch, which this workflow constrains to `main` via `branches: [main]` and which an external contributor cannot name. The point is to keep the pattern out of the file so a future edit cannot turn it into a real injection. actionlint: clean before and after. Refs: ROCM-26711 / SEC-00830
…00837) (ROCm#5110) Mythos scan finding SEC-00837 (ROCM-26712): self-hosted runners are non-ephemeral, so `docker login` credentials written by one job stay in ~/.docker/config.json and are readable by whatever runs next on that machine. aiter-test.yaml has three `Docker login` steps and no `docker logout` anywhere: build_aiter_wheels runs-on: build-only-aiter (no cleanup step at all) standard runs-on: ${{ matrix.runner }} (has "Cleanup container") multi-gpu runs-on: ${{ matrix.runner }} (has "Cleanup container") This adds `docker logout` to the two existing `Cleanup container` steps and gives build_aiter_wheels the cleanup step it was missing. All three run under `if: always()`. This is the immediate mitigation the ticket calls for, not the fix. It narrows the window but does not close it: credentials still exist on disk between login and logout, and a cancelled job may skip cleanup entirely. The actual fix is to register the runners with `--ephemeral` (or `ephemeral: true` under actions-runner-controller) so every job starts from a clean machine. That lives in the runner infrastructure, not in this repository. Existing partial mitigation, unchanged by this PR: all three `Docker login` steps are already gated on `!github.event.pull_request.head.repo.fork`, so fork PRs never write credentials in the first place. actionlint: clean. Refs: ROCM-26712 / SEC-00837
…orm_rope_cache_pts_quant_shuffle (ROCm#5027) * [ROCm][Perf] Add head_dim 512 + weightless V-norm to fused_qk_norm_rope_cache_pts_quant_shuffle Enable the fused QK-norm + RoPE + KV-cache op for Gemma4, whose full-attention layers use head_dim 512 and whose every attention layer applies a weightless v_norm (RMSNorm with has_weight=false). - rope_common.h: add warp_rms_norm_no_weight_ (RMS normalize a head with no learned gamma); apply it to V in fused_mrope_rms_kv_kernel when the new runtime flag v_norm is set, before the KV-cache write. Add case 512 to the fused_rope_rms_set_kv head_size switch (VEC_SIZE=16 at 512) and relax the head_size guard. The mrope-3D launcher is unchanged (passes v_norm=false). - fused_qk_norm_rope_cache_quant.cu / .h / rocm_ops.hpp: thread the trailing bool v_norm (default false) through the pts entrypoint and pybind. - ops/fused_qk_norm_rope_cache_quant.py: add v_norm to the python wrapper. Validated with a standalone call at head_dim 256 and 512: the 512 template instantiates/compiles, and the V-cache matches a weightless-norm reference at bf16 rounding tolerance for both widths. Co-authored-by: Claude <noreply@anthropic.com> * [ROCm][Perf][Test] Cover weightless V-norm + head_dim 512 in pts fused op Add test_fused_qk_norm_rope_cache_pts_v_norm: exercise fused_qk_norm_rope_cache_pts_quant_shuffle with v_norm on/off at head_dim 256 (Gemma4 sliding) and 512 (Gemma4 full). Asserts the V written to the cache is weightless RMS-normalized when v_norm=True and raw otherwise. Co-authored-by: Claude <noreply@anthropic.com> * [ROCm][Perf][Test] Address review: guard shuffle K write + real v_scale coverage Fix two issues from PR review: - Guard the shuffle-layout K write against silent cache corruption. That path does a single contiguous vec_t store of VEC_SIZE = head_size / WARP_SIZE elements and get_shuffle_layout_k_base() assumes they all land in one x-wide chunk (VEC_SIZE <= x). At head_size=512 / WARP_SIZE=32 that is VEC_SIZE=16, which exceeds x=8 for a bf16/fp16 cache and would corrupt K for block_size>1. Reject that config with an AITER_CHECK. An fp8 cache (x=16) still satisfies the bound, so shuffle layout at head_dim 512 with fp8 KV is unaffected. - Rework the v_norm op test to follow the file convention and add real scale coverage. The per-tensor v_scale only divides V on the fp8 quant write path -- a same-dtype cache copies V verbatim, so the previous bf16-cache test never exercised the scale. The test now uses @benchmark, is wired into __main__ with a markdown summary table, and sweeps head_dim 256/512, v_norm on/off, and (bf16, fp8@1.0, fp8@0.5) cache/scale pairs to check norm-then-quantize ordering. Co-authored-by: Claude <noreply@anthropic.com> * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: root <root@quanta-ccs-aus-k09-19.adc.amd.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…OCm#4841) * fix(topk): add acquire fence for mb radix barrier last block `radix_kernel_persistent` uses a per-row cross-block barrier in the multi-block radix top-k path. The waiting blocks observe `pass_done` with an acquire load, which also invalidates their cache state before they reload the global histogram for the next pass. The elected "last" block, however, only publishes `pass_done` with a release store and then falls through to the same plain histogram reload without ever doing an acquire/invalidate. On MI355X this can let the elected block reload stale histogram lines from the persistent workspace. If that block computes a different `local_len` / `local_k` from its peer blocks, it can leave the pass loop early while another block continues into the next barrier. The early-exiting block can then be elected in the self-reset epilogue and zero `pass_done` while its peer is still spinning on it, producing a permanent GPU wedge in the GLM-5.2 DSA indexer path. Add a `__threadfence()` plus CTA sync after the release store in the elected last-block branch. This gives the last block the missing acquire-equivalent ordering before it reloads the histogram, making both sides of the barrier observe consistent global memory before computing the next pass state. This fixes the production hang seen with GLM-5.2 TP4 + DP attention + LMCache + atomesh `dp_sticky`, where one DP rank could wedge inside `aiter::mb::radix_kernel_persistent` and then stall the whole service through DP-attention collectives while `/health` stayed green. Validation: - unpatched stock kernel soak reproduced 5 wedges in 213,400 launches / 27.3M row-launches. - fixed stock kernel soak completed 1,914,200 launches / 245.0M row-launches with 0 wedges, 8.97x the baseline exposure. - detector build changed the failure signature from `passes=[2,2,1,0]` with `STUCK` to 0 `STUCK` events over 131.4M row-launches. - end-to-end GLM-5.2 TP4 + DPA + LMCache + atomesh aiperf run has passed warmup and entered 3600s profiling without the previous hang signature. Signed-off-by: Phi-C <chenxjhit@163.com> Co-authored-by: Cursor <cursoragent@cursor.com> * perf(topk): use acquire-only mb radix barrier fence The elected last block needs device-scope cache invalidation before reloading the global histogram, but does not need release/writeback semantics. Use an agent-scope acquire fence to preserve correctness while avoiding the unnecessary release overhead of threadfence. Signed-off-by: Phi-C <chenxjhit@163.com> Co-authored-by: Cursor <cursoragent@cursor.com> * fix(topk): drain mb histogram atomics before barrier Ensure every wave completes no-return histogram atomics before block arrival, then establish agent-scope visibility after relaxed polling to prevent cross-block divergence. Co-authored-by: Cursor <cursoragent@cursor.com> --------- Signed-off-by: Phi-C <chenxjhit@163.com> Co-authored-by: Cursor <cursoragent@cursor.com>
* [CI] Mirror PR title component tags as auto-managed labels * address comments * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Xin Huang <Xin.Huang@amd.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Adds five a8w8 bpreshuffle (N,K) groups -- 1536x1536, 2048x512, 3584x7168, 7168x1024, 7168x1792 -- and three bf16 groups -- 896x7168, 7168x35840, 20480x7168. Each group covers M as every power of two from 1 to 32768 plus M=48, so a 48-row request resolves to its own entry instead of padding up to M=64. None of these had a gfx950 row before: the bf16 file carries rows for two of the three groups, but only for gfx1250. Existing rows are untouched. The tuners ran without --all, so only the newly added shapes were considered, and a key-wise comparison against the pre-tune files confirms zero modified and zero removed rows. a8w8, --libtype all -k --shape_grouped on gfx950 (cu_num=256): 85 rows, 78 FlyDSL / 4 CK / 3 CK-tile. Seventeen land on the FlyDSL 8wave pipeline and nine use splitK, all of them on 3584x7168 where K is large enough for the extra parallelism to pay off at small M. Measured against today's behaviour (no tuned row, default kernel), three runs, per-shape median of 100 iterations after 20 warmup, one GPU, with the CK and CK-tile lookup tables rebuilt from the new config first: N=1536 K=1536 322.2us -> 248.2us +22.99% N=2048 K=512 241.9us -> 175.7us +27.36% N=3584 K=7168 1797.2us -> 1373.5us +23.58% N=7168 K=1024 876.3us -> 631.4us +27.94% N=7168 K=1792 1192.0us -> 910.8us +23.59% total 4429.6us -> 3339.6us +24.61% Best single shape is M=2 N=7168 K=1792 at 16.28 -> 6.24us. Two shapes on 1536x1536 first looked like small regressions; a seven-run recheck put both sides within 0.5% of each other, which is inside the noise for a 4.5us kernel, so they are kept. bf16, csrc/gemm_a16w16/gemm_a16w16_tune.py without --with-hipblaslt, run under --compare --update_improved so a row is only written when it beats the default kernel by at least 3%. Twenty of the 51 candidate shapes cleared that bar; the other 31 are already at what the default dispatch picks and are left out. The largest win is M=1 N=896 K=7168 at 14.69 -> 6.64us (54.8%). op_tests/tuning_tests/test_config_shape_collision.py and test_csv_validation.py pass (30 tests, 37 subtests).
* Document and automate release plan * Harden AITER release automation * Update release notes after asset upload * fix: harden release automation checks * Fix manual release Docker login * Adjust release cadence anchor * Fix reusable release Docker login Signed-off-by: Xin Huang <Xin.Huang@amd.com> --------- Signed-off-by: Xin Huang <Xin.Huang@amd.com>
…ptr (ROCm#4950) * [Triton] fix(gated_delta_rule): replace removed tl.make_block_ptr for Triton 3.8 Triton 3.8 removed block pointers. tl.make_block_ptr still exists as a symbol but raises at trace time: NotImplementedError: Block pointers have been removed in favor of the tensor descriptor API so every gated_delta_rule kernel using it fails to compile. This is an API removal, not a GPU issue - it reproduces identically on gfx950 and gfx942, and is what makes op_tests/test_gdn_prepare.py fail on both MI35X and MI300X. Convert all 128 block accesses to plain pointer arithmetic with explicit bounds masks, reproducing the previous boundary_check=(0, 1) semantics: prefill/chunk_o.py 42 sites (6 kernels) prefill/fused_solve_tril_recompute.py 41 utils/solve_tril.py 35 prefill/fused_cumsum_kkt.py 10 utils/cumsum.py 2 The 2-D helper in chunk_o.py takes both strides so the transposed (K, T) views with stride (1, H * K) convert without a special case. Stores keep their fp_downcast_rounding="rtne" behaviour. Validation on gfx950 with triton 3.8.0+amd.rocm7.1.0.gitf6a045ff: op_tests/test_gdn_prepare.py 28 rows, max |err| = 0.0, all shapes / all three hidden backends (triton/flydsl/hip) * fix(gated_delta_rule): convert remaining l2norm/wy_representation block ptrs l2norm.py (4 sites) and wy_representation.py (11) still used tl.make_block_ptr, which Triton 3.8 removed. Both are on live e2e inference paths that op_tests/test_gdn_prepare.py does not exercise: gated_delta_rule.py: l2norm_fwd(q/k) when use_qk_l2norm_in_kernel=True prefill/chunk.py:109: recompute_w_u_fwd (non-fused w/u path) so a real GDN forward raises NotImplementedError at trace time. Under torch.compile this surfaces as a masked backend-compile failure. Convert both with the same plain-pointer-arithmetic pattern. Verified on gfx950: chunk_gated_delta_rule(use_qk_l2norm_in_kernel=True) now runs to finite output; l2norm_fwd matches its torch reference (max |err| 9.7e-04).
* Initial first verison of fuse_reduce_rms_mxfp4_quant_kernel(). Included changes to api call and relevant op_test. * Moved tensors descriptors for second phase into relevant section. Removed redundant layout descriptor. Removed placeholder comment. * Code Style check. * Included _mxfp4_quant_op from triton with gluon adaption. Added barrier() to sync threads. Added warning for calling gluon without proper arch. * Ruff checks
Reverts everything PR ROCm#4978 (f4e7c75) changed under `aiter/ops/triton/` back to its pre-merge state (4ad9983), plus the top-level op_tests that exercise only those Triton ops. Reverted: - aiter/ops/triton/** (159 added files removed, 12 modified files restored) * attention/{sparse_mla_dsv4_train,dsv4_indexer,mxfp8_attention}.py * SonicMoE (ops/triton/sonicmoe.py, _triton_kernels/moe/sonicmoe/, utils/sonicmoe_config_utils.py) * moe/{moe_wgrad,moe_aux_loss,moe_gemm_mxfp8,moe_gemm_per_token}.py * cross_entropy.py (+ chunked CE), quant/{quant_mxfp8,fast_transpose}.py, _triton_kernels/quant/quant_fp8_blockwise.py, gemm/basic/gemm_mxfp8.py * fusions/mhc.py DSV4 APIs (mhc_pre_dsv4 / mhc_post_dsv4 / mhc_head_dsv4, MHC_DSV4_BACKWARD_FALLBACK) and their kernels * normalization/rmsnorm.py large-M / small-N backward path * utils/_triton/arch_info.py is_cdna4() * configs/gfx942/** tune tables (125 CSV/JSON) - op_tests/triton_tests/{fusions/test_mhc.py,normalization/test_rmsnorm.py, utils/mhc_ref.py} - op_tests/{test_sparse_mla_dsv4_train,test_dsv4_indexer,test_sonicmoe, test_moe_aux_loss}.py -- these live outside triton_tests/ but import only the reverted Triton modules, so they would fail to collect if kept. Kept (non-Triton parts of ROCm#4978): - aiter/ops/fused_norm_quant_gemm.py, csrc/{include,kernels}/fused_norm_quant_gemm.* - aiter/configs/a8w8_blockscale_bpreshuffle_tuned_gemm.csv (HIP-side tune rows) - op_tests/test_topk_softmax.py (tests the ASM topk_softmax, not Triton) The reverted paths are byte-identical to 4ad9983; no commit merged after ROCm#4978 touches any of them, and nothing left in the tree references a removed module or symbol.
* [Tuning] Fill remaining Kimi-K3 bf16 GEMM shapes PR ROCm#5124 tuned 51 new bf16 shapes but --update_improved only wrote the 20 that beat the default by the 3% threshold, leaving the rest to fall back to default dispatch. The tuner had picked a best backend for all 51; the other 31 were ties rather than losses. A/B those 31 against default dispatch on a single gfx950, 3 rounds of 100 iters after 20 warmup, and add the 26 that are not slower: N=896 K=7168 M=1024,4096,16384,32768 N=7168 K=35840 M=1,2,4,8,32,48,256,1024,2048,16384,32768 N=20480 K=7168 M=1,2,8,32,64,128,256,512,1024,4096,4096 Largest gains are at N=20480 K=7168 small M (+12~16% over default). 5 shapes measured slower than default and are left untuned. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * [Tuning] Group Kimi-K3 bf16 tuned rows by (gfx, N, K) The new shapes from ROCm#5124 and this PR were appended at the tail, so each (N, K) group ended up split across several places in the file. Gather the rows of every group this round touched into one contiguous block sorted by M, at the position where the group first appears. Row contents are unchanged - reordering only. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * [Tuning] Add the three torch-winner bf16 rows back For M=16/128/512 at N=7168 K=35840 the tuner's best candidate is torch, which is also what default dispatch runs, so the A/B measured the same kernel on both sides and the sub-1% deltas were noise. Record them. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * [Tuning] Pin the last two bf16 shapes to torch M=8192 at N=896 K=7168 and M=48 at N=20480 K=7168 measured 2.2% and 3.9% slower with the tuner's pick (triton / flydsl) than with default dispatch, so record what default already runs - torch - rather than a slower kernel. Completes the 51 new shapes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
…m#4787) * [gfx950] Optimize Minimax M3 scoring & top-k kernels * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - formatting change * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - formatting change * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - fix compile issue * Fix msa scoring kernel test * [HIP][JIT][gfx950] Minimax-M3 - Fix module name for scoring & topk kernels --------- Co-authored-by: root <root@smci355-ccs-aus-n07-09.cs-aus.dcgpu>
…code Interface (ROCm#5011) * [FlyDSL] Add variable-length decode TopK Add a compile-time-specialized multi-CTA radix-select path with optional deterministic output for the existing per-row decode interface.
situv2_and_mul_quant computes the SiTUv2 gated activation and the per-token FP8 quantization of its result in one pass, so callers read the [m, 2*d] bf16 input once and write only the fp8 output plus the scale. Two forms are dispatched. When the row fits one chunk per thread the activations stay in registers and a single DPP wave reduction produces the scale; a single-wave block skips LDS and both barriers. Wider rows stage through FP32 LDS and sweep the row in a loop. The block/vector-width ladder is tuned per d on gfx950 so the block stays close to d / VecSize -- a loose fit leaves whole waves idle. At m = 32768 this reaches 5698 GB/s at d = 2048 and 5938 at d = 4096, the measured pure-copy bandwidth ceiling. On a TP8 decode trace the kernel replaces a 154.4 us activation plus the 146.8 us MLP share of the quant kernel with one 140.3 us kernel per step, 2.15x on the replaced set. That set is ~1% of total GPU time, so end-to-end stays within run-to-run variance. Output is within 1 ULP of the fp32 reference. The fp8 bound comes from opus::finfo<fp8_t>::max(), so both the OCP e4m3fn range on gfx950 and the fnuz range on gfx942 are handled. Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
* docs: update README * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
* Widen prefill N-tiles for Kimi-K3 * Drop the MoE sort/tile M-block from 32 to 16 * Clean up comments * Clean comment * Rename config file to a16w4 instead of fp4 * Default the stage 1 waves_per_eu to None so the int4 port is unaffected The int4 stage 1 registry does not emit a waves_per_eu key, so _flydsl_stage1_wrapper's numeric default of 3 was the value that reached the a16w port. That was harmless while the port hardcoded waves_per_eu=None, but now that the value is forwarded it clamps occupancy on a path this work does not touch: all 88 tuned kimik2_i4 rows were measured with no such attribute, and rocdl.waves_per_eu=3 also appends _w3 to the JIT cache key, so every one of them would take a cold rebuild and run a differently scheduled kernel. Default to None instead, which is what stage 2 already does. Every other stage 1 family sets waves_per_eu explicitly in the registry, so the default only applies to int4 and the fp4 families keep the axis live. * Point ROCm#2 Replace the copied timings on the retiled prefill rows with measured ones The eight token>=2048 rows changed tile_n but kept the us1/us2/us/tflops/bw values from the pre-retile file, so they described a kernel that is no longer the one named. That is worse than cosmetic: update_config_files sorts by us and keeps the lowest per shape when merging tuned CSVs, so those numbers can decide which row wins. Re-measured each row at its own parsed config, block_m, tile_n, tile_k, b_nt, xcd_swizzle, k_wave and waves_per_eu taken from its own kernelName, median of 3 repeats at 200 iterations. us is us1+us2 and tflops/bw are recomputed with gemm_moe_tune.py's own formulas, which reproduce every unchanged row in this file exactly. The new values run 15 to 19 percent above the replaced ones because this host is capped at 1000 W against a 1400 W TBP, the same host the block_m=16 rows were measured on. No other tuned CSV covers these eight shape keys, so the merge dedup has nothing to compare them against. err1 and err2 are left as they are. tile_n does not change gemm1's K-reduction order, so its output is bitwise identical across 64, 128, 192 and 384, and err1 genuinely does not move. gemm2 changes only the order of its bf16 atomic scatter, which cannot shift a value printed to one decimal place. Dispatch is untouched: block_m, ksplit, kernelName1, kernelName2, run_1stage, xbf16 and flat are byte-identical for all 32 shape keys. * Make the TILE_N wave-partition asserts test what their messages say Both asserts checked the per-wave column count after a floor division, so a TILE_N of 64k+1 to 64k+3 slipped through: 193 // 4 is 48 and 48 % 16 == 0, so the gemm2 assert passed for a TILE_N that is not a multiple of 64, which is exactly what its message claimed to enforce. A reader trusting the message would either reject 192, which is legal, or accept 193, which is not. Test the divisibility directly instead. gemm2 wants TILE_N % 64 == 0; gemm1 partitions 4 // k_wave N-waves, so it wants TILE_N % (16 * (4 // k_wave)) == 0, which is 64 at k_wave=1, 32 at 2 and 16 at 4. No registered tile_n changes verdict: of every a16w4 and a16wi4 stage 1 name that clears the pre-existing >= 16 assert, zero pass the old condition and fail the new one, and the stage 2 set of 128 and 256 is unaffected. Verified reachable with D_INTER = 193 * 128, where the divisibility and 256-alignment asserts both pass and TILE_N=193 now fails the wave-partition check. * Rename the untuned twin to match its renamed tuned file The tuned config became kimik3_a16w4_tuned_fmoe.csv but its untuned input kept the old kimik3_fp4_ prefix, leaving the only unpaired fmoe config in the directory; every other pair shares a prefix. That is a trap rather than a cosmetic issue. The orphan lists exactly the 32 shape keys the tuned file covers, and the tuner derives its output name from the input name, so a retune driven off it writes kimik3_fp4_tuned_fmoe.csv and recreates the deleted file beside the new one with all 32 keys duplicated. The next config merge then hits the duplicate-shape check and refuses to start. No runtime effect either way: both loaders glob and exclude any name containing untuned. test_csv_validation.py checks a fixed seven-file list under configs/, not model_configs/, so CI would not have caught this. * Let the a16w4 tuner reach the configs the tuned CSV ships Three filters in gen_flydsl_2stages_task made the shipped kimik3 a16w4 rows unreproducible, so a retune would silently regenerate different ones. blockM was gated to 32, 64 and 128, so no block_m=16 candidate was ever timed even though the a16w4 and a8w4 stage-1 registries emit t16 names and the int4 sibling already allows 16. Any retune of 3584/384 would have discarded the decode configs this work is built around. Added 16; the s1_tile_m != blockM filter below still drops the dtypes with no t16 names, so a4w4 gains nothing and a8w4 gains the 32 candidates its registry already advertises. Stage 2 refused tile_n=256 for bf16 outright, which is why no run could have produced the eight t32x256x128_atomic rows. The comment was right that the hazard is LDS over-allocation at large tile_m, but the code did not implement that. The port requests tile_m*(tile_k*2 + tile_n*4), so check the bound directly. Confirmed at both sides of it by forcing real codegen, since compile_gemm2_a16w4_port only builds the launcher: tile_m=128/tile_n=256/ tile_k=128 is 163840 bytes and runs, while tile_k=256 is 196608 and fails with "local memory (196608) exceeds limit (163840)" -- the predicted figure exactly. The shipped tile_m=32 config is 40 KiB. The num_acc_n filter used a floor divide, so it would pass a tile_n like 96 or 160 to a compile worker that then hits the kernel assert mid-sweep. Made it the same divisibility test the kernel now uses. The tuner would now generate the config named by all 32 rows of the CSV, up from 21. The v2 path's identical blockM gate is left alone: it returns early unless the activation is fp4 or fp8, so a16w4 never reaches it. * Correct two registry comments that overstated the rules they describe The k_wave comment said the option set applies to "tile_m in {16,32}", but the guard is tm == 32 or (tm == 16 and is_a16w4), so fp8xfp4 gets k_wave only at tile_m=32 even though its registry does emit 32 t16 names, none of them with a _kw suffix. Someone adding a8w4 k_wave tuning would assume t16 _kw names exist and hit "Invalid FlyDSL kernel name". The int4 comment still said kw=1 "requires tn >= 64" after the filter moved to a divisibility test. Under that test tn of 80, 96, 112 and 160 are all >= 64 and silently dropped, so anyone extending tile_ns past 16, 32, 64, 128 would read the wrong rule and debug a name the registry never emitted. Comments only; the parsed AST is unchanged. * Scope the tile_m=16 speedup claim to what was actually measured The comment attributed the win to E=896 generally, which reads as covering both tuned inter_dim families. Only inter=384 has block_m=16 rows; inter=512 was never swept for it and all 16 of its rows stay at 32. It also gave no basis for the figure, so a reader comparing it against the us column of this CSV lands on 1.02x to 1.65x instead, because those rows were measured on two different hosts. State the measurement instead: 1.18-1.35x on the isolated GEMM pair at E=896 inter=384 for token<=512, which is the same-host A/B both this branch and the original sweep produced, and note that inter=512 is excluded. Comment only; the parsed AST is unchanged. * Reduce the added comments to one line each Eight of the comments this branch adds ran to two, three or four lines. Each is now a single line, and the two that sat inside pre-existing multi-line comments are back to the original wording with only the one clause that the code change invalidated rewritten, so the diff shows one added line rather than the whole paragraph. The whole branch now adds 14 comment lines across five files, none adjacent, longest 88 characters. Comments only; the parsed AST of every file is unchanged. --------- Co-authored-by: root <root@smci350-rck-g03-d09-31.rck.dcgpu>
* CI: add extended test workflow * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * CI: avoid extended test reruns on PR edits --------- Co-authored-by: Leo <drleonid@amd.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
FlyDSL is required after ROCm#5116, so the removed helper breaks test collection and silently disables decode dispatch. Co-authored-by: Cursor <cursoragent@cursor.com>
…des ROCm#4870) (ROCm#5142) * feat(skills): add validate-kernel-pr, a runtime validation layer for kernel PRs review-pr is static: it never builds and never runs. Three failure modes are invisible to it, and a 14-PR backtest of review-pr alone found 8 of 17 known defects, so the gap is real rather than theoretical. This skill produces validation_report.json as the evidence base: merge_sim -> gpu_claim -> runtime_compat -> test_policy -> correctness -> index_width_scan -> verdict Two stages exist because a green test suite is not evidence: - correctness runs the PR's own tests AND a shape grid this skill owns (non-toy, boundary/odd, long-context). Seeding a dropped tail guard into a FlyDSL softmax kernel leaves the repo's own suite green -- its non-aligned shapes are commented out, only one aligned shape runs -- while the grid at N=2000 / N=257 fails it. - test_policy compares tolerances head-vs-base before running anything. A change that widens a tolerance while leaving the kernel path unchanged makes the suite unable to fail, and pytest still reports green. Verified against three seeded mutants plus a clean baseline: baseline passes every stage; dropped tail guard blocks via the grid alone; loosened tolerance blocks via test_policy alone; wrong index unit blocks via the repo tests. The report states what it did not do, so it cannot overclaim by omission: arch_coverage marks each arch runtime or compile-only (a gfx950 host cannot validate a gfx942 claim), isolation records the real level, and degraded_mode marks COMPILE_ONLY when no GPU was claimable. runtime_compat exists because a pinned prebuilt runtime drifts behind the tree and the resulting ImportError otherwise reads as a defect in the PR. scan_index_width.py ships alongside it; see the review-pr D9 change. * fix(skills): make review-pr's D9 trigger structural, and require validation evidence Four changes, each from a measured failure in a 14-PR backtest of this skill. 1. D9 never fires. Its trigger was a list of variable names (token_id, seq_start, batch_offset, total_tokens). Three real 32-bit overflow defects used none of them -- stride_out_batch, block_id, physical_block, context_kv_idx -- and int32/int64/overflow appears zero times in the whole verdict card for aiter#1674, which carries two of them. The trigger is now structural: an index-shaped value multiplied by a stride-shaped value on a line with no 64-bit widening. scan_index_width.py is the mechanical pre-filter and flags all three; a bare model missed the same family, so this is not something to leave to recall. 2. REPO was hard-coded to ROCm/aiter, so the skill could not review a FlyDSL PR at all. It now takes a repository argument, or owner/repo#N. 3. FlyDSL/Triton kernel PRs now require a validation_report.json and load validate-kernel-pr. A kernel PR's own green suite is not evidence. 4. The verdict card caps at 5 findings, ordered most-severe first, and states its validation evidence. This skill averaged 7.1 findings per PR across the backtest; dropping everything judged noise still left 5.4, so the cap has to be explicit. Without a report the card says [static-only review], and no finding may then assert runtime behaviour as fact. Backtest numbers for the frozen version being changed here (blob 3fdb11af, 14 PRs, 17 provenance-checked labels): 76 of 100 findings useful, 0 factually wrong, 8/17 defects found. * docs(skills): correct validate-kernel-pr's documented interface The Invocation section showed `validate_pr.sh --pr 4394`, a flag the script does not implement -- it exits 2 with "unknown arg --pr". That is the hallucinated-API failure review-pr Step 6 exists to catch, shipped in the skill meant to catch it. Document the interface that is actually implemented and verified: the caller supplies the worktree and the pytest target. Every flag named in the doc is now present in the argument parser. Also declare what is absent rather than leaving it to be inferred: - no PR fetch orchestration; choosing the right --tests target from a diff is the unsolved part, and a wrong target produces a confident green - perf and claims stages are reserved in report_schema.json and emitted by nothing, so a report today carries no performance evidence and a review must not read a missing perf stage as "no regression" * fix(skills): run the D9 index-width scan in Step 1, not mid-checklist A scan that Step 5 asks for does not happen. In a controlled 14-PR run of the revised D9 text, no session invoked the scanner and the arm caught 0 of 3 known overflow defects -- the same as the unrevised skill. Moving the invocation into Step 1's fetch block does make it run: the session transcript for aiter#1674 shows the scanner executing and its candidate list entering the context. D9 now works that list instead of asking for a scan. This is necessary but NOT sufficient, and the PR text says so: with the scan in Step 1 the arm still caught 0 of 3. What closes the gap is supplying the production-scale facts the 5xx-gate demands -- see the limitations section. * fix(skills): attribute test failures against a baseline control Found by running validate-kernel-pr end to end on a real PR for the first time (FlyDSL#969, conv3d fp8 padded-M store guard). Both the pre-fix and post-fix trees failed the same test shape -- one unrelated to the PR -- and the harness reported "the PR's own test target fails" as a blocker in both. It charged pre-existing red to the author, which is the misattribution this skill exists to prevent. With --patch, the suite now runs on the base first and a failure is only a blocker when the base is clean; otherwise it is a note saying the target is red with and without the change. The same run goes from BLOCK to PASS with the pre-existing failure recorded rather than blamed. Also carries the production-scale table as a separate file printed by Step 1 beneath the scan candidates. Its effect is documented in the PR description and is small: 1 of 3 on the target defect family, versus 0 of 3 without it. * feat(skills): add downstream-impact-check, built from a real cross-repo incident Downstream tests here are label-gated and skipped by default, so the question "can this break vLLM or SGLang" decides whether they run -- and it is normally answered from memory. aiter#4530 is why that is not good enough. It changed one file, aiter/ops/triton/_triton_kernels/moe/moe_routing/topk.py: kernel internals, no public signature touched. It broke vLLM's gpt-oss MXFP4 MoE on MI355 after the 0.1.16.post5 -> 0.1.19 bump (vllm#49361, test plan blank), needed a hotfix that names it as the companion fix (vllm#50859), and SGLang reverted its aiter pin the same week (sglang#32879). scan_downstream_consumers.py walks aiter's imports upward from the changed files and stops at the first module a downstream checkout imports. On #4530 it reconstructs the chain from the internal topk kernel through moe_routing and moe_op_gemm_a16w4 to vLLM's aiter_mxfp4_w4a8_moe.py, and correctly reports no reachability for SGLang, which does not import that subtree. The first version asked "did a public symbol change" and answered "no downstream impact" for that exact PR. Signature is the wrong question; reachability is the right one. That earlier criterion is recorded in the file so it is not reintroduced. Validated on one positive and one negative, which is thin -- the skill says so, along with the rest of what it cannot see: static imports only, five hops, the downstream's current main rather than the pinned version, and no ATOM checkout. Wired into review-pr Step 1 alongside the other scanners, because a check the reviewer has to remember to run does not run. * fix(skills): prevent false kernel validation clearance Separate advisory review output from deterministic validation, make every missing evidence path explicit, and preserve exact base/head attribution. Commit the pinned mutant corpus so these safeguards remain reproducible. * fix(skills): reject unexercised validation paths Treat shadowed FlyDSL source, unused shape-grid hooks, and base-run artifacts as inconclusive instead of allowing them to produce false validation clearance. * fix(skills): harden validation evidence boundaries Require an exercised shape-grid handshake, isolated base/head state, and runtime-backed architecture coverage so incomplete evidence cannot be accepted as a clean result. * fix(skills): bind reports to the live base Resolve the current base branch tip instead of trusting the historical PR merge base, so validation reports become stale as soon as their merge simulation does. * fix(skills): require observed execution evidence Prevent green pytest from becoming clearance unless validator-owned instrumentation observes the expected route and shapes, JUnit records real executions, and runtime identity is captured. Align report semantics with process exits and reject untrusted FlyDSL runtime rebuild claims. * fix(skills): support script validation targets Run script targets through their real entry points so successful non-pytest validation stays inconclusive instead of producing a false BLOCK. Ship the AMD GPU picker so the validator can discover an idle translated HIP device. * Treat GPU activity as optional and hand the worktree back clean amdsmi_get_gpu_activity was a hard dependency at three call sites, so a host where that single query fails degrades the whole run to NO_GPU even with idle GPUs available. @zufayu hit this on MI308X / ROCm 7.0, where the call returns AMDSMI_STATUS_UNEXPECTED_DATA (43) while enumeration, BDF, ASIC and VRAM queries all work. PICKER could not route around it because the executor re-queries the API inline rather than going through the picker's selection. Activity is now optional. A new read_activity() helper is shared by the picker and both inline probes, and gpu_claim.idleness_basis names the evidence the claim rests on: activity+vram, or vram-only when only resident VRAM separated the devices. Unknown stays distinct from zero -- the previous `gfx if isinstance(gfx, int) else 0` substitution reported an unavailable metric as a measured idle GPU. The gpu_claim skip also stops conflating "GPUs present but none idle", an environment fact, with "AMD SMI is unqueryable", a portability gap in the validator. The import fallback did not probe $ROCM_PATH/share/amd_smi, where the bindings ship on ROCm >= 7.1, so on those containers the picker exited 2 before reaching any activity query -- the same NO_GPU symptom from an unrelated cause. Separately, merge_sim applied the patch but cleanup was guarded by BASE_ACTIVE, which is set only inside the baseline path. A run that skipped correctness exited with the patch applied, and when the baseline path did run, cleanup's restore_head re-applied it, so every path left the caller's worktree patched and the next run reported not isolated-clean against the caller. The application now carries its own flag and is reverted on exit, including on interrupt. Reported by @zufayu on ROCm/aiter#4870. * Let script targets carry grid and receipt evidence A target driven by a `__main__` guard could reach neither correctness_s1_grid nor execution_receipt, so INCONCLUSIVE was the hard ceiling for every aiter `op_tests/*.py` -- the repository's standard test shape. Neither skip described a property of the target: * the route profiler is sys.setprofile plus a receipt write, and pytest was only where the hook happened to be installed; * the grid was injected solely through an environment variable, while these targets take shapes on the command line. That is a limit of the injector. Reported as skips, both read as honest abstention while actually hiding a capability gap -- the failure mode this skill exists to prevent. The probe is now installed for script targets by run_script_with_probe.py, which executes the file under runpy with run_name="__main__". It calls the probe's own hooks rather than re-implementing the profiling, so `producer` stays truthful and the tested PR still cannot forge a receipt. The probe module is generated whenever a route is named, not only on the pytest branch; without that the runner had nothing to import. --shape-arg names the target's own CLI flag for the grid. The flag is named by the caller rather than guessed, because a wrong guess appends argv the target silently ignores. The hook is held to the same standard of proof as the env-var channel: the flag must appear in an add_argument call, and the existing invalid-grid probe still has to make the target fail before the stage is credited. A receipt is now validated whenever a route was named, including when no grid was configured or the grid channel could not be established. Two branches used to abandon the receipt along with the grid, discarding evidence already collected. With no grid it asserts route execution and nothing about shapes, which is all it is then entitled to claim. report_schema.json relaxes the PASS constraint `test_selection.runner: const pytest` to `enum [pytest, script]`, since a script target can now supply both missing stages. shape_arg is declared but deliberately not required, so reports from earlier validators stay valid. Verified on ROCm/aiter#4538 (gfx950, MI355 OAM), base 78b9440e/891788643: * script target with --shape-arg -s: PASS 9/9, exit 0, review-pr consumable, receipt naming _auto_variant over the four injected production shapes * unaccepted --shape-arg flag: correctness_s1_grid skip, receipt still pass * seeded defect (fused -inf fill dropping [e, seq_len_kv)): BLOCK, exit 1, with correctness_s1_grid failing independently of the PR's own suite * tests/test_validator.py 24 passed; PASS/INCONCLUSIVE/BLOCK reports and one pre-shape_arg report all validate against the schema * Ask the target whether it needs a GPU, and name what that answer rests on An unclaimable GPU suppressed all four correctness stages, so a target that needs no device reported nothing rather than reporting what it could establish. The target is now asked directly: when no GPU was claimed it is run once with no visible device, and test_selection.gpu_requirement becomes not-required only if it passes AND executes at least one test. The executed count is what makes this evidence rather than a guess -- a suite guarded by skipif(not torch.cuda.is_available()) also exits 0 having proved nothing, and a verified fixture covers that case. This is deliberately an observation, not a judgement about the diff. A Python-level dispatch change reroutes kernels without touching kernel source, and ROCm/aiter#5089 decides whether 34 gfx950 kernels compile from a seven-line helper, so no static rule over changed paths could settle it. Asking the target cannot be wrong in the dangerous direction: a target that needs a device fails or skips without one and stays required. PASS is deliberately NOT relaxed. `complete` still requires gpu_claim: pass, so this path cannot produce a clearance that was previously unreachable -- verified by a case where merge_sim, runtime_compat, test_policy, baseline_control, correctness_repo_tests, correctness_s1_grid, execution_receipt and index_width_scan all pass and the verdict is still INCONCLUSIVE with arch_coverage empty. gpu_claim stays `skip` because no device was claimed, which remains the fact; only the suppression is lifted. Two notes on scope, because this is a contract change rather than a defect fix. It edits behaviour @zufayu explicitly did not ask to change ("the gate is right and INCONCLUSIVE is the correct output"), and it changes an invariant the suite asserted, so test_no_gpu_is_inconclusive_and_every_skip_is_declared was updated and test_no_gpu_withholds_correctness_from_a_target_that_needs_a_device added to keep the original protection under test. It also does not by itself let a shape-less bugfix reach PASS: correctness_s1_grid has no passing path without a consumed grid. Also splits the last conflation in the gpu_claim skip. A host with no GPUs now exits 3 with "this host reports no GPUs" instead of borrowing the activity-unavailable wording, which blamed the validator for an absent device. Co-authored-by: Cursor <cursoragent@cursor.com> * State the promotion bar in review-pr's header @zufayu asked for this in item 4 of ROCm/aiter#4870: the bar belongs in the header because a header is read on every use while an issue sinks. #5128 was opened for the measurement work and closed as premature, so without this the condition for ever leaving advisory tier is recorded nowhere. Half of it was already present -- "its judgement is stochastic and never blocks a merge". What was missing is the bar itself: the tool stays advisory until false clearance is measured and near zero for the families that raise a red verdict. That number is named as the one that matters, distinct from recall or a spot check, and stated not to exist today because no committed replay corpus establishes it. Ownership is attached to whoever proposes a rule edit or proposes gating, which is where the trigger actually is. Also completes the item-4 headline ask. The report was already optional in the code (`VALIDATION_REPORT="${3:-}"`, the no-report branch, "absence of a report is not itself a blocker", and Step 8's NOT RUN line), so nothing gated on it. The one place that read as mandatory was the front-matter description, which said to invoke with "an explicit validation report path"; it now says a review without one is a supported outcome. Takes the second option offered for exempt target classes and names them on the FlyDSL kernel row, because the combination is what item 4 objected to: a CPU-only target claims no GPU and therefore no architecture, and a bugfix with no shape dimension has no grid for correctness_s1_grid to consume. Both are structurally unable to reach PASS, so their INCONCLUSIVE is the expected output and that author must not be asked for a passing report. Co-authored-by: Cursor <cursoragent@cursor.com> * Let review-pr accept a script target's PASS `Let script targets carry grid and receipt evidence` relaxed the PASS-requires-pytest assumption in the validator and in report_schema.json, but review-pr re-derives the verdict independently and still required `selection["runner"] == "pytest"`. A script target that legitimately reached PASS was therefore rejected by the consumer as self-contradictory: validation verdict contradicts its stages/findings: expected INCONCLUSIVE, got PASS Two of the three places that encode the contract were updated and the third was not, so the change silently made the evidence unusable at exactly the point it was meant to be consumed. The coverage-basis check immediately above already handled `script`, which is what made the omission easy to miss. Verified against the report from ROCm/aiter#4538 on gfx950 (script target, --shape-arg -s, PASS 9/9): rejected before, "validation report accepted for head 22850902..." after. tests/test_validator.py 25 passed. * Have review-pr triage validation need, and run it when it can A judgement the checklist asks for mid-review is a judgement that does not get made, so Step 1 now classifies the changed paths itself and, when the PR changes runtime surface and ships exactly one test target, invokes bin/validate-kernel-pr and consumes the result through the unchanged identity gate. Producing a fresh head-bound report is not the same act as adopting one found on disk, so "never auto-load ./validation_report.json" still holds. The verdict line gains the two states it was missing. A README fix used to report the same "NOT RUN" as an unvalidated kernel rewrite, which teaches a reader to skip the line; "N/A - no runtime surface changed" and "required but not run" are different facts and now read as such. Co-authored-by: Cursor <cursoragent@cursor.com> * Call the validator from the skill, not from bin/ The auto-run reached outside .claude/skills for bin/validate-kernel-pr, whose only addition over validate_pr.sh is a PR-number front end: it re-fetches the diff and re-resolves the base tip. Step 1 already holds both, so it now creates a detached worktree at the base it recorded and hands that, its own pr.diff and the head OID straight to the validator. Deriving those a second time was not merely redundant. main can advance between the two gh api calls, and the resulting report names a base the review's own identity gate then rejects as stale -- a self-inflicted failure on a report we just produced. Handing over the recorded base makes repo.base agree by construction, since the validator reads it as rev-parse HEAD of that worktree. The dependency also now sits beside the index scanner and the report schema, and fails loudly like they do instead of warning and continuing, since a skill tree missing one of the three is broken rather than differently configured. Finally, each way the run can give up records why. Triage's blocker only covers attempts never made, so a run that was attempted and failed printed "required but not run: None" -- the exact uninformative verdict line this step exists to remove. Co-authored-by: Cursor <cursoragent@cursor.com> * fix(skills): make the D9 index-width scan an AST pass with no name lists The scanner matched a regex over raw diff text whose operand character class contains neither a comma nor a space, so it could not span a broadcast subscript -- and ptr[:, None] * stride is the standard Triton pointer-arithmetic idiom. Measured on aiter#4978, the PR that introduced the moe_wgrad overflow later fixed by #5132: both real defect lines were missed, the one moe_wgrad candidate emitted was an already-int64 line, and 390 candidates were produced overall. The defect lived in the syntax tree; the scanner was reading text. It also matched operand names against two hard-coded lists. Both are gone. A multiplication is now a candidate when it feeds an addition chain, at least one operand derives from a non-constexpr parameter of the enclosing kernel, and no operand is widened -- counting a widening carried in from an earlier line, which is how #5132 fixes it. Post images come from --source-root, else the diff's own index blobs; a file whose post image cannot be recovered is reported as NOT SCANNED, never dropped, because an empty candidate list that silently excluded files is not evidence of absence. On #4978 it now reports 4 of 4. On #5132 it reports none. review-pr's D9 text claimed a structural trigger while still defining index- and stride-shaped by name; it now describes what the scanner actually does. * feat(skills): deliver the S1 shape grid through pytest parametrization Applying this validator to four open aiter FlyDSL kernel PRs (#5167, #5141, #4992, #5011) reached INCONCLUSIVE on all four, for one reason: correctness_s1_grid skipping is on its own enough to force it, and the stage is inert on this repository's dominant test shape. Zero of the seven files in op_tests/flydsl_tests/ read a shape env var or parse a shape CLI flag; every one declares its shapes as literals inside @pytest.mark.parametrize. --shape-argnames adds parametrization itself as a third channel, via a generated plugin whose pytest_generate_tests substitutes the grid. The burden of proof is unchanged: a deliberately unusable grid must still make the target fail, so a target that ignores the grid produces a skip and never a pass. The invalid-grid row keeps its arity and poisons its values, so the probe fails inside the target rather than inside the plugin. A mark binding the requested names together with others is refused rather than stripped: removing it would leave those others as unfilled fixtures and the target would error for a reason unrelated to the kernel. Two channel defects found alongside it. The env probe overwrote the CLI probe unconditionally, so supplying both flags discarded a working CLI channel and then reported the env channel's absence as the reason no grid ran; the probes are now independent and run_pytest dispatches on the channel that actually probed positive. And the skip text read 'kernel exposes no configured shape override' even when the cause was the validator ignoring the named channel -- grid_channel_reason now names each channel tried, what was found, and what the target does offer instead. * fix(skills): stop reporting pass on evidence that was requested and never collected Three stages decided their status from 'this code path completed' rather than 'evidence was actually collected', and published pass on empty fields. execution_receipt: the profiler sampled f_locals only on the call event, so it saw parameters and nothing else. A route deriving its shapes in its body -- the common case for a kernel whose arguments are tensors -- could never satisfy --shape-vars, and the receipt reported pass with executed_shapes: []. Observed on three separate PRs. The frame is now sampled again on return, keyed by the frame OBJECT rather than id(), whose values are reused once a frame is freed; the first version of this fix used id() and dropped the second of two sequential calls. When capture was requested and produced nothing, the receipt now carries a shape_capture block saying the receipt asserts route execution and makes no claim about shapes. test_policy: the tolerance pattern matched only literal atol=/rtol=, so a file declaring DEFAULT_REL_TOL = 2e-2 reported zero tolerances and the check passed on an empty list. Loosening a named constant is exactly the m2 mutant this stage exists to catch, and it was undetectable. Named constants are now captured, and a tolerance passed by name is recorded so an empty list is never read as 'none'. * fix(skills): claim a GPU that is actually free, and name the basis honestly PICKER resolved from PATH before the picker this skill ships. The shipped picker is what prints idleness-basis:, so a PATH copy that omits it won by default and the report published idleness_basis: 'unknown' beside a concrete gfx_activity_before_pct: 0 -- an unavailable reading presented as a measured idle one, which SKILL.md argues against in its own text. Resolution is now explicit PICKER, then shipped, then PATH. The picker is deterministic and returned one device, and a contended flock -n gave up. Concurrent validators therefore all selected the same GPU and all but one reported no idle GPU while seven sat idle; observed twice in one run with every lock free between attempts. The picker gained --all, and the executor walks the eligible ranking until a lock is taken. * test(skills): cover the validator changes Thirteen tests for the AST scanner, the three grid channels, receipt honesty and the GPU claim. Verified against the pre-fix sources: 5 of 6 scanner tests and all 8 of the executor tests fail there, each for the intended reason. The exception is the earlier-line-widening test, which passes pre-fix coincidentally -- the old regex also failed to span token_row[:, None], for its own wrong reason. test_json_count_is_deduplicated is rewritten rather than restored: it asserted the name-list contract that was deliberately removed. * fix(skills): stop the validator from publishing its own faults as the PR's Round-2 application to five real FlyDSL kernel PRs found six ways a fact about the run was published as a defect of the code under test. Three were introduced by the pytest grid channel itself. --shape-argnames with exactly ONE name passed argnames as a list while unwrapping rows to scalars. pytest sets force_tuple only for a string argnames, so collection died with "object of type 'int' has no len()" and the executor published that crash as 'the PR adds this target and its independent shape grid fails' -- a BLOCK verdict against the author, on three separate PRs. One name is the dominant shape in the targets this channel was built for. The partial-overlap guard was evaluated over the whole FILE, so an unrelated test whose parametrize mark overlapped the requested names disabled the channel for every test in it. The plugin decides per metafunc; the probe now decides per test function, as it always should have. A relative --patch was resolved against the caller's cwd in one place and the worktree in another, so merge_sim reported 'patch does not apply to the current base' and returned BLOCK. Paths are resolved once, before anything moves. head-repo and head-grid shared one receipt path through their phase directory, so a grid run that crashed during collection erased a receipt that had already proved the route, and the report then said the route never ran. Receipts are per label, and one helper decides which speaks for the head run: a phase that observed nothing never speaks over one that observed something. The invalid-grid probe treated any non-zero exit as proof the grid was consumed. On a held-out PR whose module could not be imported, and again when the plugin itself crashed, the channel was credited although no shape reached the kernel. The control run must now pass before the probe's failure means anything. JUnit errors were counted as executed tests, so a collection error credited arch_coverage with runtime coverage on the strength of the error itself. WIDEN_ATTRS was matched case-sensitively; FlyDSL writes fx.Int64(...), so explicitly widened FlyDSL code was reported as an overflow candidate. Found only on the held-out PR -- the four PRs used while fixing all happen to write tl.int64. Two long-standing gaps closed alongside them. --tol-table was parsed, validated and compared against nothing; it now yields a finding when the suite accepts error beyond the loosest reference supplied. And runtime_identity published native_artifacts: [] as a measurement when the probe runs before the target and imports only aiter -- it now states what that emptiness does and does not mean. The pytest channel delivers the grid as test PARAMETERS while the receipt records the ROUTE's locals. Requiring containment across those two vocabularies produced a missing-shapes skip on a run whose every grid case passed; the requirement is recorded, the cross-namespace assertion is not made, and review-pr's ingestion gate knows the difference. * test(skills): cover the false-attribution fixes Twelve tests for the second batch. All twelve fail against the pre-fix sources and none pass there, so each one discriminates. The single-name --shape-argnames test asserts the injected values reach the route rather than the target's own literal, which is what the false BLOCK hid. One needed the caller's directory nested two levels rather than one: with a single level, a relative --patch resolves to the same file whether it is read from the caller's cwd or from the worktree, and the defect is invisible. The fixture's synthetic HIP index moved from 7 to 57. The validator flocks /tmp/gpu-<index>.lock, and an unrelated job on this host holds the lock for 7, so every GPU-claiming test degraded to NO_GPU for a reason that had nothing to do with the code. No assertion was weakened; only the synthetic device number changed. * fix(skills): gate the pytest grid on parametrize VALUES, not just argnames Round 3 found the false BLOCK on #5141 was not gone. The plugin crash was fixed; the crash had moved into the author's file, which makes the misattribution harder to see rather than less real. The gate read a parametrize mark's argnames and never its values. A target parametrizing a single `case: dict` therefore passed it, the validator substituted integers, and the target raised `TypeError: 'int' object is not subscriptable` -- published as '[blocker] the PR adds this target and its independent shape grid fails', against an author whose own 138 tests passed in the same report. SKILL.md already documented that such targets stay INCONCLUSIVE; the executor did not implement it. Both the AST gate and the plugin now require the mark's own values to be scalar cells, and values of unknown shape count as not scalar -- a grid this channel cannot express must cost an INCONCLUSIVE, never a blocker aimed at an author. Three more false attributions closed with it. review-pr asserted `executed == tests - skipped` while the validator had been corrected to subtract errors. They disagree only when errors > 0, which is the shape of a report carrying a real runtime blocker -- so a report that had correctly found one was rejected at ingestion and could not be used. Worse than the overclaim it replaced, and caught only by re-running the case that had produced a true BLOCK. The invalid-grid probe's causality guard had been added on the base side only, so the head branch still credited the channel when the target failed for an unrelated reason. And review-pr turned a grid the caller supplied into a requirement even when no channel had carried it, so the receipt's honest empty list read as a contradiction and discarded exactly the runs carrying the accurate diagnostic. The scanner claimed kernel scope in its docstring and its class name while visiting every Python function: on a held-out PR all four candidates were host-side float FLOP accounting and none were in the 916-line kernel. Device scope is now a predicate -- a constexpr-annotated parameter, or a jit/kernel decorator -- and host-scope hits are listed separately rather than dropped, because that predicate is a spelling test and a wrong answer should cost a glance, not a miss. The index-width scan now reads post images from the applied worktree, so modified files are examined rather than disclosed as unscanned. gpu_requirement no longer reports a conservative default in language that reads as an observation. * fix(skills): scope the scan to the enclosing kernel, and reject a caller's typo at the door Two defects found by re-running the cases that had produced decisive verdicts. A grid whose rows did not match --shape-argnames produced a BLOCK. The arity check lived inside the plugin generator, whose exit status run_pytest never read: the generator failed, the PREVIOUS phase's plugin file survived on disk, head-grid re-ran the invalid-grid sentinel it still carried, and that failure was published as 'the PR adds this target and its independent shape grid fails'. A check whose result nobody reads is worse than no check -- it creates the appearance of a guard while the system falls back to stale state. The arity check now runs at argument parsing, where a caller's mistake cannot become a finding about the PR, and the generated plugin is removed before regeneration so no phase can inherit another's. The index-width scan was classifying real kernel code as host-side. _scan_body used ast.walk, which descends through nested defs, so every expression inside a nested helper was scanned in the ENCLOSING function's scope: its parameters, its widened locals, its device/host verdict. Four candidates inside a @flyc.kernel body were demoted to host scope, and a miss matters more here than noise. The walk now stops at function boundaries, and scope, parameter provenance and widening follow the scope chain -- a nested helper closes over its kernel's parameters, so it inherits them. Fixing that exposed a second bug it had been masking: visit_FunctionDef recursed with generic_visit, which visits a statement's CHILDREN, so a nested def that IS a statement of the body was never dispatched at all. Evidence unchanged where it should be and corrected where it was wrong: aiter#4978 still reports 4 of 4 on the moe_wgrad overflow, #5132 still reports none, #5141's nested-kernel arithmetic moves from 4 host to 5 device, and the held-out #5025's host-side FLOP accounting stays out of the device list. Two abstentions also stated the wrong cause -- 'does not take all of these as test parameters' for a target that does take them, when the real refusal was non-scalar mark values. The probe now returns the refusal it actually made. * fix(skills): let auto-validation reach the pytest shape channel The auto-validation block added on this branch forwards REVIEW_SHAPE_ENV and REVIEW_SHAPE_ARG but not the parametrization channel, so the one channel that reaches this repository's dominant test shape was unreachable from the caller that was just wired up. None of the seven files in op_tests/flydsl_tests/ reads a shape env var or parses a shape flag; all declare their shapes as literals in @pytest.mark.parametrize. Also records where a caller learns which channel a target actually offers, so a wrong guess costs one run rather than a reading of the target's source. * Measure a kernel PR's cost, not just its correctness A kernel PR could pass every stage while running slower, because nothing timed it. Add a `perf` stage: bench the target once with the patch reversed and once with it applied, both inside the base phase's lock on the same GPU, and compare. The comparator (scrape_perf.py) reads both of aiter's timing formats -- the `df.to_markdown()` tables from `@benchmark`/`run_perftest`, and the older `[perf] ... ck avg: N us` lines that 11 kernel targets print. A scraper that only knew tables would burn a full base+head run on those and report nothing. Three choices worth recording, each forced by a measurement rather than by taste: - Take the *minimum* over per-column medians, not the mean. A PR touching the ck path leaves `torch avg` at 1.0, and a mean would let that untouched reference column bury the regression next to it. - Run each side three times and keep the best sample per cell. Five warm runs of unchanged test_layernorm2d.py gave `torch avg` 14.24 14.64 14.18 14.23 14.31 (1.03x spread) but `ck avg` 13.10 20.98 20.70 13.28 13.17 -- 1.60x, and bimodal rather than noisy. Single-shot at a 0.95 threshold would have called unchanged code a regression about half the time. Best-of-3 brings the same data to 1.014x, worst ratio 0.986. - Snapshot the worktree around the timing run and roll back only the paths it dirtied. Without this a bench that writes a results file (see op_tests/test_gemm_a8w8.py, test_moe_2stage.py) fails the base phase's cleanliness check, which sets BASE_READY=0 and silently skips head correctness entirely. Measured: the same target went PASS with --no-perf and INCONCLUSIVE with perf on. A perf stage that quietly disables correctness is worse than no perf stage. A regression under the threshold writes a should-fix finding, so the deterministic verdict becomes NEEDS_WORK and the process exits 1. review-pr's identity gate now cross-checks the stage's status against its own numbers, so neither a laundered pass nor a laundered fail survives, and its P6 rule fires only when no usable stages.perf exists. The harness detector is shared in spirit by both skills and pinned by a test, since a silent disagreement there means one skill demands a measurement the other cannot supply. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * Describe the validator this branch actually ships The perf stage landed six commits ago and its own skill still says it does not exist: "the schema reserves both -- and the script emits neither. A report today carries no performance evidence, and a review must not read the absence of a `perf` stage as 'no regression'." That instructs a reader to disbelieve a stage that now measures, gates, and can turn a verdict into NEEDS_WORK. review-pr says the opposite in three places, so the two skills contradicted each other on the one stage a reviewer is most likely to act on. Document `perf` where the implementation lives: the flag table gains --perf-args and --no-perf, the env knobs gain PERF_TIMEOUT, PERF_REPEAT, PERF_THRESHOLD and PERF_MIN_ROWS -- all of them ${VAR:-default} overrides with their own argument validation, so they were public already and only undocumented -- and the stage list gains a section that says why best-of-N is what makes a 0.95 threshold usable and why every non-measurement path reports `skip`. The verdict section now states that PASS does not imply a timing comparison happened. What remains genuinely unimplemented is narrower than the old bullet claimed: reproducing the specific numbers a PR description states. That is now what "Not implemented yet" says. Three smaller corrections in the same vein: - "review-pr never builds and never runs" has been false since review-pr started invoking this script itself. The split is at judgement, not at invocation, and the text now says so. - The comment describing bin/validate-kernel-pr as an existing wrapper outlived the wrapper; bin/ is not in the repository. The reasoning it carried -- do not re-resolve a base that main can advance past -- is worth keeping, so it stays, phrased about a front end that does not exist rather than one that does. - The mutant manifest's tolerance_table disagreed with the file it describes: test_softmax.py at the pinned base uses f16=1e-2 and bf16=2e-2, not 2e-3 and 1e-2. Inert today, because reference tolerances are recorded and not compared, which is exactly how a wrong number survives unnoticed until something starts reading it. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * Stop the report schema promising evidence nothing produces report_schema.json is the contract review-pr validates against, so a field declared there reads as a shipped capability. Three did not hold. `claims` declared a stage that reproduces the numbers in a PR description, with an `unreproducible` array. No producer writes it -- not validate_pr.sh, not validate_evidence.py, not scrape_perf.py. Its only other mention was an allowlist entry saying it may be absent. It is removed rather than left standing; the honesty rule that promised `[unreproducible]` tagging now describes what `stages.perf` actually carries. `source_sha` was hardcoded null while the same honesty rules listed source SHA as recorded evidence. It is fillable: the resolved module has a path, and that path is usually inside a checkout. Now it records that checkout's commit, suffixed `-dirty` when the tree has uncommitted changes -- which the head phase always does, since it runs with the candidate patch applied, so the bare commit would name the base and describe a tree that is not the one under test. A module resolved from an installed package stays null: attributing a commit to a prebuilt binary is precisely the provenance claim this skill refuses to make. `reference_tolerances` is recorded and never compared, so --tol-table influences no verdict. Wiring it up is not possible as the data stands -- the table is keyed by dtype while the scraped tolerances are a positional list, with nothing to join them on -- so the field now says that it is context and not a gate, rather than implying otherwise by sitting next to the comparison that is one. Also collapses the PASS gate's two byte-identical correctness blocks into definitions/passing_correctness_stage, and has its `stats` $ref the execution_stats definition that was already there instead of respelling it inline. 74 duplicated lines in the clause that decides what earns a clearance, where divergence between the copies would change the answer silently. Verified equivalent: a clean PASS report still validates, and executed=0, failures=1, errors=1, exit=1, status=skip, missing stats and a stats object missing `skipped` are each still rejected on both stages. Adds the fields the producers already emit but the schema never described -- gpu_claim.post_run_note and .requirement_note, perf.regressed_rows[].base_repeats and .head_repeats -- which passed only because additionalProperties defaults to true. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * Name the channel the shape grid actually travelled on --shape-arg was added so script targets could receive an S1 grid, but the code around it still assumed the grid could only be an environment variable. Three consequences, one of them visible in a report. GRID_HOOK_OK was assigned twice. The --shape-arg AST scan set it, then an unconditional second block re-derived it from the --shape-env scan. Supplying both flags therefore proved the env var while run_pytest put the shapes on argv -- the report attested a channel the run did not use. Only-one-flag runs worked, which is why this survived. The channel is now resolved once, before either scan, and --shape-arg wins for a script target because that is what run_pytest does. A --shape-arg naming a flag the target does not accept fell through to the branch that reports "no shape grid was configured". Both cases skip, so the reason is the only thing distinguishing a validator that could not find the hook from a caller who never asked for one -- and the wrong one of those is an environment fact standing in for a capability gap, which is the overclaim-by-omission this skill exists to prevent. The branch conditions now test whether a grid was requested at all, so the same input reports hook-not-found, names the flag, and records test_selection.grid_channel. --shape-arg on a pytest target gets its own answer rather than the same silence: argv never reaches it. The runtime skips likewise said "shape environment variable" whichever channel was in play. They name the real one now. run_pytest took "$SHAPE_ENV=$GRID" and re-split on the first `=`. For a CLI-only run that worked only because an unset SHAPE_ENV left a leading `=` the split then removed, so the shapes rode inside a string shaped like the channel they were not using. It takes the grid value directly and picks the channel from the same decision that credited the hook. Collapses the two grid-less branches, which were a 17-line receipt block duplicated verbatim with two different skip strings. The new test fails on the previous script for the reason that matters -- 'no configured shape override' unexpectedly found -- not merely on the field added to carry it, which is why the note assertions come before the grid_channel one. 40 pass. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * Find the index-width overflow the D9 scan was written for The scanner matched a regex over raw diff text whose operand character class contains neither a comma nor a space, so it could not span a broadcast subscript -- and `ptr[:, None] * stride` is the standard Triton pointer-arithmetic idiom. It also matched operand names against two hard-coded lists. Measured on aiter#4978, the PR that introduced the moe_wgrad overflow later fixed by #5132: 390 candidates emitted, none of them the four real defect lines, and the single moe_wgrad candidate was `off_expert * stride_dwe` -- a line whose operand is widened with `.to(tl.int64)` forty rows earlier. Recall 0 of 4, on the defect family the rule is named for. The defect lived in the syntax tree; the scanner was reading text. Rewritten as an AST pass, taken from #5174. Both name lists are gone. A multiplication is a candidate when it feeds an addition chain, at least one operand derives from a non-constexpr parameter of the enclosing kernel, and no operand is widened -- counting a widening hoisted to an earlier line, which is how #5132 fixes it. Scope follows the scope chain rather than ast.walk, so a nested helper inherits its kernel's parameters instead of being scanned in the wrong frame. Post images come from --source-root or the diff's index blobs; a file whose post image cannot be recovered is reported as NOT SCANNED rather than dropped. It reports 4 of 4 on #4978 and 0 on #5132. Two defects fixed on top of what #5174 ships. AugAssign was never scanned. _scan_body matched ast.BinOp with an Add op, but `a_ptrs += BLOCK_K * stride_ak` is an ast.AugAssign, which is not a BinOp -- so the most common Triton pointer-advance idiom was a blind spot. The constexpr filter in _is_compile_time was written for exactly that expression and could never reach it. Adding it moves #4978 from 296 candidates to 306; the moe_wgrad four and the #5132 zero are unchanged. WIDEN_CALL_RE was left behind by the regex implementation this replaces and had no remaining reader. The --json path returned `0 if not unscanned else 0`. It stays 0, because both callers require that: #5174's own test asserts the plain and --json runs exit 0, and validate_pr.sh turns a non-zero exit into a skipped stage rather than a louder finding. Making the exit status honest needs validate_pr.sh to pass --source-root or to read the `unscanned` field, so the reason is recorded where the return is instead of a ternary whose two branches were the same value. Base's IndexScannerTests asserted the name-list contract that was deliberately removed, down to a JSON key that no longer exists; #5174's rewritten IndexScannerTests and ScannerScopeTests replace it. 48 pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> * Stop the receipt reporting pass on shapes it never captured The probe sampled a frame's f_locals on the `call` event only, so it saw parameters and nothing else. A route that derives its shapes in its body -- the common case for a kernel whose arguments are tensors -- could never satisfy --shape-vars, and the receipt published `pass` beside `executed_shapes: []`. Deciding status from "this code path completed" rather than "the evidence was collected" is the failure this skill argues against in its own text, and it was in the probe. From #5174: the frame is sampled again on return, keyed by the frame OBJECT rather than id(), whose values are reused once a frame is freed and would drop the second of two sequential calls to the same route. When capture was requested and produced nothing, the receipt now carries a shape_capture block saying it attests route execution and makes no claim about shapes, so no consumer can read the empty list as evidence that no shapes were needed. The evidence checker stops counting errors as executed tests. Errors are collection and fixture failures where the body never ran, so a target that failed to import was crediting arch_coverage with runtime coverage on the strength of the collection error itself. This also makes the checker agree with review-pr's ingestion gate, which already asserts `executed == tests - skipped - errors`: they disagree only when errors > 0, which is the shape of a report carrying a real runtime blocker, so the two had to be corrected together or a correct report would be rejected at the door. runtime_identity stops publishing `native_artifacts: []` as a measurement. The probe runs before the target and imports only the module under test, so an empty list describes that import's reach and not the kernel's; the basis string now says which. validate_receipt gains a `--grid-channel` argument, used to suppress the cross-namespace containment assertion when the grid travelled as pytest parameters. It defaults to empty and validate_pr.sh does not pass it on this branch, so that branch is inert until the pytest grid channel lands. 50 pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> * Let the GPU picker rank every eligible device, not just one The picker returned a single device and the executor took a contended `flock -n` as a refusal. Concurrent validators therefore all selected the same GPU and all but one reported "no idle GPU" with seven sitting idle -- observed twice in one run with every lock free between attempts. A deterministic picker and a non-blocking lock are only safe together if the caller can fall through to the next candidate. From #5174: --all prints the eligible devices in ranking order rather than the top one. The executor side of this, walking that ranking until a lock is taken, lives in validate_pr.sh and is not in this commit, so nothing consumes --all yet; the default single-device output is unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> * fix(skills): stop the validator reporting coverage it did not obtain Applying validate-kernel-pr to ROCm/aiter#4538 (a rewritten gfx950 FP8 MQA logits kernel) returned PASS in 93 seconds. The PR has a reproducible correctness failure -- flydsl_fp8_mqa_logits asserts at num_heads=16 on gfx950, where the pre-PR base computes it correctly -- and its central performance claim went unmeasured. Every gap below was found by running the tool, and each fix owns a generic invariant rather than that PR's shape. A shape grid that duplicates the target's own defaults is not a control. All three requested cells were already in the target's `--shapes` default, so the "independent" grid re-ran a strict subset of correctness_repo_tests and was credited `pass` -- the exact duplication SKILL.md says the stage exists to prevent. Grid cells are now compared against the flag's declared default; `test_selection.grid_independence` reports the outcome, and a passing duplicate is downgraded to `skip`. A duplicate that FAILS keeps its `fail`. A coverage axis the shape flag cannot reach could not be requested at all. `--grid` is one ordered tuple on one channel. A target whose head counts, dtypes or window modes live on separate flags had entire configurations unreachable however the grid was spelled. `--axis NAME=FLAG:v1;v2` adds them, under the same burden of proof as `--shape-arg`: the flag must be declared in the target's own argparse AND observed refusing a deliberately invalid value. An axis that fails either is dropped and NAMED, never dropped quietly. A script target's `executed: 1` stood for 56 graded cases. It is the same number a silently-returning target produces, and it earned the same `arch_coverage: runtime` on basis `script-exit-zero-with-output` -- a statement about the process, not the kernel. `stats.observed_work` now carries route calls counted in that run's own receipt, and a script run that observed no work earns no architecture credit. head-repo and head-grid shared one execution receipt. The second erased the first, and with the grid cells a subset of the defaults a receipt written by either run satisfied `--grid`. One receipt per run; `execution_receipt.receipt_scope` names the one that was published. "The PR adds this target, so base has nothing to time against" is not true. The file being new does not make the code it drives new. When the target only exercises an entry point that exists on base, the target is transplanted into the base tree and times the pre-PR implementation through the same harness. Because that spans two trees, `--perf-control-column` names a column the patch does not touch and the stage skips unless it reproduces within PERF_CONTROL_TOL. On #4538 this yields 56 matched rows, the kernel column 1.29x faster on head, and the untouched Triton control within 0.0%. scrape_perf matched 0 of 56 comparable rows. A header with no recognised unit was treated as identifying, but aiter bench tables print unlabeled MEASUREMENTS (`flydsl rel`, `triton err`, `speedup`) that differ between base and head by construction. The strict key is still tried first; only if it matches nothing is a relaxed key used, dropping identity columns whose cells are non-integral. `row_key_basis` records which. The target ran with the reviewer's environment attached. `env VAR=... cmd` adds to the inherited environment, so unmerged third-party code could read every token in the calling shell -- reproduced, with real credentials reaching a stage log. The target now runs under `env -i` plus a name/prefix allowlist and a secret-shaped denylist, and `isolation.target_environment` reports what was passed through. The exit code was read back out of `--out`. That made a previous run's report a fallback source of truth: a run killed before finish_report exited on the earlier run's verdict. `--out` is removed at startup and the code comes from a verdict file inside this run's own $WORK. Ten regression tests, each verified to fail against a8d46dcb0 and pass here; the suite is 62 and green. The synthetic picker moved off a real device index, which was making the suite contend with genuine validations on the same host. * fix(skills): let the mutant replay reach its own verdict check The driver reverse-applied each case's patch unconditionally after the run. But validate_pr.sh already hands the worktree back with the patch reversed, so that second reverse-apply failed with "patch does not apply" on EVERY case, and under `set -euo pipefail` it aborted the driver before the summary block that compares each verdict against the manifest. The mutants were being discriminated correctly the whole time -- m0 PASS, m1 and m3 BLOCK in correctness, m2 BLOCK in test_policy -- but the suite could not say so, and its exit status was 1 regardless of the result. A regression asset that cannot report itself green is not one. Reverse only when `git apply -R --check` says the patch is still applied. End to end on ROCm/FlyDSL@421935cc with the installed flydsl 0.3.1 (verified identical to that commit's python/flydsl tree), the replay now prints "mutant replay matched all expected verdicts". Found by running the committed replay as a held-out case, not by reading it. * fix(skills): say which fact a skipped comparison rests on Three defects that only appeared on ROCm/aiter#5172, a fresh held-out PR frozen before any of this work started and never looked at while fixing. A requested axis vanished from the report when it could not be honoured. #5172's target is collected by pytest, so argv-borne axes cannot reach it and axis_state became "unusable" - correct - but test_selection.axes was published as [], so a reader could not see that an axis had been requested at all. An empty list beside a non-"none" state loses the request itself, which is exactly the silently narrowed test space these fields exist to make visible. Axes are now recorded as asked for, with hook_proof "not-evaluated" when the run never got far enough to scan the target for the flag. grid_independence published a false statement about the target. Its default reason, "the channel exposes no declared defaults to compare against", was emitted whenever the comparison did not happen - including when the channel had been demoted for an unrelated reason. On #5172 that is untrue: the target's -c flag does declare a default list, and all three requested cells were outside it. The reason now names which of the several skip paths applied. "Red on both baseline and head" is an attribution, not an explanation. #5172's target defines pytest nodes AND parses argv in its module body, so pytest imports it at collection with its own argv and argparse exits. The file is green run as a script - the perf stage in the SAME run proves it, three times per side. The report said the target was red and gave no hint that the runner selection was the cause. test_selection.runner_risk now names that structurally, and a run that executes nothing under the selected runner cites it. Thirteen new regression tests in all, each verified to fail against a8d46dcb0; the suite is 63 and green. * fix(skills): stop charging a runner-selection artefact to the PR author A second script-target case, ROCm/aiter#5081, was picked after the first held-out round showed that the mechanisms added here had only ever executed their positive path on ROCm/aiter#4538. Selection was structural and outcome-blind ("an aiter PR that ADDS an op_tests script target with a shape flag and a sibling axis flag"); it immediately produced a false blocker. "Defines a test* function" is not "pytest can collect it". aiter's dominant op_tests convention is a SCRIPT whose worker happens to be named test_<op>(m, d, dtype) and is called from main() with real arguments. pytest collects it, cannot supply the parameters, errors -- and the validator published "the PR adds this test target and it fails on head" against an author whose target is green as a script with the very shapes the run had requested (verified by hand: -m 97 513 -n 512, all checkAllclose passed). A test* function is now only counted as a pytest node when it takes no required positional parameters or carries a parametrize/fixture decorator. #5081 selects `script`, the grid and the axis reach it, and no blocker is raised. `--runner` lets a caller settle the case the classifier still gets wrong, recording both the forced choice and what selection had said. A route behind functools.wraps could never be observed. A frame's identity is f_globals["__name__"] + ":" + f_code.co_name, and wraps copies __name__ onto the wrapper object while leaving co_name alone -- so aiter's whole @compile_ops family executes as `aiter.jit.core:wrapper` and naming the op matched nothing. The probe now resolves the declared route to code objects through the __wrapped__ chain; the string match remains as a fallback for a module that cannot be imported. On #5081 the receipt goes from observed '' to a real observed route, and INCONCLUSIVE now rests on the honest remaining gap (a (*args, **kwargs) dispatch wrapper binds no shape locals, so the receipt reports missing shapes rather than passing). Two fields answered the same question differently. When a proven axis rescues a duplicate shape grid, the override reached stages.correctness_s1_grid.independence but not test_selection, which had been written earlier -- so one report said "duplicates-target-defaults" and "adds-coverage" at once. Both are now written from the same decision. A control-gated perf stage still published the numbers it had rejected. `status: skip` beside a median_ratio and a regressed_rows list reads as a regression that was merely not acted on. When the control column moves outside PERF_CONTROL_TOL the ratio fields are dropped. Verified on #5081 end to end: novel grid -> grid_independence adds-coverage with axis `-n` proven (target_defaults 384/768/1024/4224, novel 512/1536); duplicate grid + no axis -> the stage downgrades to skip and both fields agree; duplicate grid + proven axis -> rescued, and they still agree. Seventeen new regression tests, all verified to fail against a8d46dcb0; the suite is 68 and green. --------- Co-authored-by: Jin Pan <jin.pan@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: zhiding512 <zhiding512@users.noreply.github.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
… fallback (ROCm#4994) * fix(fmoe): fuse stage-1 fp8 quant on the heuristic FlyDSL fallback get_2stage_cfgs' tuned path selects the fused "_fp8" stage-1 variant per shape, which folds stage 2's mxfp8 quant and scale-sort into stage 1's CShuffle epilogue: stage 1 then writes one byte per element instead of two and the standalone quant+sort kernel disappears. Most rows in tuned_fmoe.csv carry that suffix. The heuristic fallback never considered it, so any shape without a tuned row ran the separate conversion even though the fused variant existed for every tile. DeepSeek-V4 at EP8 takes the fallback, so it paid this on every MoE layer -- measured at 116us per layer, 8.8% of decode GPU busy time. Fused unconditionally: a sweep of inter_dim 384..3072 found no crossover, with end-to-end fused_moe faster fused in 14 of 16 cells (-0.4% to -4.2%) and the two exceptions (+1.0%, +0.1%) inside run-to-run spread. The tuned table's preference for unfused at inter_dim=384 reflects its joint stage1+stage2 choice, not a standalone loss: timing stage 1 alone charges the fused arm for its epilogue while never charging the unfused arm for the quant kernel that fusion deletes. AITER_S1_FUSE_FP8Q=0 disables it and AITER_S1_FUSE_FP8Q_MIN_INTER_DIM reimposes a floor. * test(fmoe): force the fallback without a tuned CSV and clear the cfg cache Two defects in the test added with the previous commit, both of which made it report success without checking anything. It forced the heuristic fallback with a header-only config_file, which only exists on newer aiter, so the test could not run at all against an older tree. Emptying the module-level cfg_2stages cache reaches the same branch and works either way. More seriously, get_2stage_cfgs is lru_cached on its arguments while the env knobs are not arguments, so after the first case every later case got the first one's metadata back. That made the two cases expecting fused pass vacuously and the two expecting unfused fail. Clearing the cache per test makes all four discriminate: the same shape now returns "fp8" or "" depending only on the knob under test. * fix(fmoe): collapse S1 fp8-quant fusion to one clear env switch Address review feedback: replace the two-env scheme (AITER_S1_FUSE_FP8Q + AITER_S1_FUSE_FP8Q_MIN_INTER_DIM) with a single on/off toggle AITER_MOE_FUSE_STAGE1_FP8_QUANT (default on; =0 keeps the separate quant). Drop the inter_dim floor -- the sweep found no crossover -- and trim the redundant comment block now that the env name is self-explanatory. Behavior is unchanged by default. Tests updated to the single switch (floor tests removed); both pass on gfx950. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…copy) (ROCm#5202) * make _fold_seqlen_indptr cudagraph safe Signed-off-by: Micah Williamson <micah.williamson@amd.com> * [Claude] add unit tests Signed-off-by: Micah Williamson <micah.williamson@amd.com> * [Claude] fix formatting and tests Signed-off-by: Micah Williamson <micah.williamson@amd.com> --------- Signed-off-by: Micah Williamson <micah.williamson@amd.com>
* [gfx950] Optimize Minimax M3 scoring & top-k kernels * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - formatting change * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - formatting change * [HIP][Kernel] Minimax-M3 scoring & topk kernel changes - fix compile issue * Fix msa scoring kernel test * [HIP][JIT][gfx950] Minimax-M3 - Fix module name for scoring & topk kernels * Add index skipping, static FP8 scaling, and packed LBHNC support so vLLM can use one AITER kernel while preserving existing cache contracts. * Use origin/main composable_kernel (15e12dd7). Co-authored-by: Cursor <cursoragent@cursor.com> * Emit unit-scale e4m3 index_q from fused QK-norm so 4787's MiniMax-M3 score kernels can consume it. * update format * update CK version * fix the CI test --------- Co-authored-by: root <root@smci355-ccs-aus-n07-09.cs-aus.dcgpu> Co-authored-by: ukannika <uma.kannikanti@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: root <root@smci355-ccs-aus-m11-29.cs-aus.dcgpu> Co-authored-by: Chuan (Richard) Li <chuali@amd.com>
Two files, both new, on top of main: op_tests/bench_gfx1250_combo.py combined bench op_tests/triton_tests/attention/test_mla_v4_triton.py MLA v4 triton reference The bench imports the top-level @benchmark sweep fns from the op_tests siblings and runs each over its own shape axes, printing only the per-op summary tables -- all the JIT/ROCTracer/import noise is silenced at the fd level. Two suites: --perf for hardware-oriented single-op numbers, --dsv4 for the DeepSeek-V4 shapes. Token sweeps come from AITER_BENCH_TOKENS; ops whose axis means something else, or whose usable range is fixed, pin their sweep in the source and say why. Everything else the branch used to carry has reached main on its own since: test_pa_sparse_prefill.py (ROCm#4926), test_mega_moe_gfx1250.py (ROCm#5052), test_flydsl_qk_norm_rope_quant.py and the mla_v4 hsa artifacts. Only these two files were ever unique to it, so this is what is left to upstream. Notes worth carrying, all measured on gfx1250 / 20260827-28: a16w16 does not pre-check the >4 GiB operand limit. That guard belongs to one fallback path -- opus_dispatch_a16w16_gfx1250 searches the tuned table first and only reaches check_shape_4g on a miss, en route to the split-K kid whose launcher builds the 32-bit gmem descriptors. A tuned 4wave_wl_co winner never gets there. Predicting it in Python skipped shapes that tuning had already made runnable, so the kernel is left to raise and the exception is recorded as a row. It also checks the error ratio the UT returns against _A16W16_MAX_ERR: all four M=65536 shapes come back err=0.96-0.99 while every other row is 0 or ~1e-5, and nothing in the UT raises or warns, so those used to print as data. a8w8_blockscale sweeps M from 1024. Below that, get_CKGEMM_config's M -> get_padded_m -> nextPow2 retry lands on ROCm#4773's M=16/M=64 gluon rows, and the UT's extra "ck strided x_scale" check (line 120, added by ROCm#4406 and gated on ck_preshuffle alone) hands triton a stride != 1 specialization that fails to compile in make_llir. The mxfp8_128 path declares its layout with is_x_scale_transposed=True and never reads the stride, so that check tests nothing there. Fixing the gate is upstream's call; starting at 1024 keeps this bench clear of it. Verified 36/36, err=0, 2207-7003 TFLOPS. mla_v4_prefill is pinned to n=1024 and mla_v4_prefill_fp8 drops nnz_prefix=8192; both are kernel/verify faults, measured and documented at the pins. Co-authored-by: Yu <jiaolyu@amd.com>
Three token constants were hardcoded and never consulted the variable --
_INVERSE_ROPE_TOKENS, _MLA_DECODE_TOKENS, _MEGA_MOE_TOKENS -- with comments
saying so ("Pinned, not env-driven", "deliberately not consulted"). The
reasoning was that a global token count means the wrong thing for those
ops, which is true of the default but not of an explicit request: if a
caller sets the variable, that is their decision to make, including asking
for a shape the op is known to fail on.
All nine token constants now go through _tokens(): unset, each op runs its
own default and says at its constant why that default is not the shared
list; set, the variable wins everywhere and the file does not argue with
it. Behaviour with the variable unset is unchanged.
The module docstring said "Two ops ignore it and pin their sweep in the
source" -- no longer true, and it only listed two of the five ops whose
default differs. It now lists all five with the reason for each.
Should have been part of ROCm#5076; amended in after that PR was pushed.
Co-authored-by: Yu <jiaolyu@amd.com>
mori_ep measured only the bf16 wire, which is not the leg DSv4 serves on.
The receiver hands the dispatched payload straight to the expert GEMM as its
A operand, and that GEMM is a4w4 -- ATOM's serve script pins MEGA_WIRE=fp4
with AITER_FORCE_A8W4=0 -- so bf16 alone measures a path production does not
take, and misses the wire whose payload is a quarter as wide.
Sweeps bf16 and fp4 by default, one child each: bench_ep.py reads $DISP once
at import and builds the transport for that dtype, so the tiers cannot share
a process. $DISP still overrides, now comma-separated, and is forwarded
unvalidated -- mori owns the value set, this file does not restate it.
fp4 rows are labelled UNCHECKED. mori forces its own comparison off on fp4
("fp4 combine is too lossy to compare"), so CHECK=1 is silently dropped
there; without the label a passing fp4 row reads as verified next to a bf16
row that actually was.
Co-authored-by: Yu <jiaolyu@amd.com>
) ROCm#5084 rewrote test_pa_sparse_prefill.py and mla_v4_prefill stopped running against it: --nnz-prefix is gone, so the child aborted with "unrecognized arguments" before reaching a kernel. - drop --nnz-prefix. The nnz axis no longer exists: the CSR is generated from --mode under --seed (sparse draws a random nnz per row, dense fills every row), so nnz is an outcome, not an input - delete the mla_v4_prefill_fp8 op. Its entire content was that nnz sweep, and it has no equivalent under the new UT. mla_v4_prefill already covers fp8 -- and now covers the backends too, since ROCm#5084 compares opus/asm/triton - _space_table: decide a data row by counting numeric fields instead of testing the first one. The new table leads with prec/mode (bf16, dense), so the first-field test dropped the whole table and the op reported "no result rows" while the UT itself had exited 0 Raise _MLA_PREFILL_TOKENS to 1024..16384. Re-measured on the new UT (b45-2, one process per tier): 1024 through 16384 all clean, where 4096/8192/16384 faulted on the old one. 16384 is the DSv4 prefill chunk and the reason this op exists, so this is the coverage the pin was costing us. 65536 stays out. It faults with a memory access at 0x7f2ddbec0000 and writes an 89 GB coredump doing it -- a third of the free disk on a shared host -- and it is past the chunk size the model prefills anyway. Both the tier sweep and this op run --no-verify, hardcoded at the call site, and that is now load-bearing rather than a speed choice: with verification on, even n=1024 dies at the first case (fp8/dense, fault at 0x43000), so the reference or the comparison is what breaks, not the kernel under test. These are therefore timings from an unverified kernel. The comment says so, because a16w16's M=65536 rows looked exactly this healthy until _A16W16_MAX_ERR was added and caught err=0.99. Co-authored-by: Yu <jiaolyu@amd.com>
) chip_info runs rocminfo twice per process -- once for the arch, once for the CU count -- and rocminfo takes a per-device rocm_smi mutex on its way in. One process is fine, and the nine single-GPU ops never noticed. A torchrun op starts four ranks at once, and they contend for that mutex. Both multi-GPU ops have now lost a run to it. On b45-1 a rank lost the mutex and aborted: init_mutex /rocm_smi_renderD128: unlock timed lock, ret: 1 terminate called after throwing an instance of 'std::runtime_error' what(): Allgather operation failed The allgather is the symptom, not the cause: the rank died first and took the collective with it. On b45-2 the other shape of the same problem -- four rocminfo processes sat in the mutex for minutes, one wedged in D state, the op never produced a line, and even docker stop hung waiting for the driver to let go. Detect once here, where the call is serial, and hand the answer down: GPU_ARCHS -> get_gfx_list() skips _detect_native() CU_NUM -> get_cu_num_custom_op() skips its own rocminfo Both are read from the environment before either shells out, so a child that inherits them runs no rocminfo at all. Set on os.environ in main() for the children that inherit our environment, and setdefault() in _run_child for the ones handed an explicit env -- setdefault throughout, so an exported value from the caller still wins. The value is what this process detected on this machine, not a hardcoded gfx1250, so it is the same answer the child would have computed. Verified in the child's /proc/<pid>/environ: GPU_ARCHS=gfx1250, AITER_GPU_ARCHS=gfx1250, CU_NUM=256, matching get_gfx()/get_cu_num() here. mega_moe on b45-1 after the change: rc=0, all 36 configs, zero mutex warnings, and zero rocminfo processes observed while it ran. Before it, the same op on the same host was rc=1 with 6 mutex warnings. Co-authored-by: Yu <jiaolyu@amd.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist