fix(GFX1100-TG200): skinny_gemm OOB write + MoE silu gate rounding

[ghazni101]

Row: GFX1100-TG200
Issue: #2889
Depends on: #2807 (T25)

Summary

One bug fix:

1. skinny_gemm OOB write: C[m + y + n * M] can write past M when m + y >= M (partial tile at the tail of the M-strided loop). Add if (m + y < static_cast<uint32_t>(M)) guard before the store.

The MoE silu gate rounding fix (originally item 2) has been split to #2957 on row/BACKEND-ROCM-MOE-SILU-ROUNDING with its own committed spec and BACKEND-ROCM ownership.

Benchmark

A/B interleaved, 5 pairs, Qwen3.5-4B Q4_K_M, 256 tokens, temp 0, seed 0:

Pair	A (T25 chain)	B (T25+Bugfix)	Delta
1	32.475	32.413	-0.19%
2	32.389	32.327	-0.19%
3	32.403	32.387	-0.05%
4	32.380	32.391	+0.03%
5	32.422	32.381	-0.13%
Median	32.403	32.387	-0.11%

Noise — correctness fix, not a performance lever.

Token identity

PASS — identical output to T25 chain.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM-5-2 [OMP]

[VikashLoomba]

Current-main reconciliation at c796fea41f74fe90b8cf78190eeb2a5b4c977449: these changes need separate owning flows before a port.

1. The skinny store is not a reachable overflow on current main's production dispatch. rocm_skinny_gemm.hip:26 fixes kYtile=2, and the sole production WvSplitKBT caller at rocm_matmul_hipblaslt.hip:577-582 requires N > 8 && N % 2 == 0. Commit f38c1edc4 landed that guard with feat(rocm): wvSplitK skinny GEMM for decode M<=4 — the #487 decode GEMM lever #506. The committed rocm-skinny-gemm.md spec explicitly records this defect class and its sentinel mutation. test_backend_cross_device.cpp:2377 covers odd features, and lines 2418-2442 check the guard band. The leaf commit's {tokens=2,K=256,features=254,YT=4} failure describes an experimental tile setting, not current main's YT=2. Keep that regression with any future tile-size or eligibility change. Do not add a current-main bug fix for a production path whose precondition already excludes the case.
2. The ROCm MoeSiluMulK at rocm_moe_router.hip:32-38 still omits the gate-dtype intermediate narrowing. This is an independent BACKEND-ROCM activation flow. Existing test_backend_cross_device: MoeSiluMul bf16 exactness reds on gfx1200 ROCm too, the ROCm counterpart of #1802 #1954 records the ROCm exactness failure and is listed under Owed in rocm-kquant-nwarps-decode.md. Main's CPU implementation narrows at cpu_ops.cpp:733. Active vLLM e126687a9 defines the same intermediate cast in csrc/libtorch_stable/activation_kernels.cu::silu_kernel, then multiplies in compute and packed_compute. This source difference supports a focused repair candidate, but it does not prove that narrowing resolves every previously reported mismatch. Port the active upstream fixtures and run the pinned operation before acceptance.

The original issue #2889 and this cumulative PR still carry the campaign row. Reduce the activation repair onto current main with explicit BACKEND-ROCM ownership and its own committed spec and review. Preserve the conditional skinny regression in its actual tile/dispatch flow. Record-anchor changes and deleted tests are not part of either kernel repair.

This is a source and history audit. No new GPU run, accepted numerical result, or performance claim is made.

[localai-org-maint-bot]

Held back with the rest of the GFX1100-TG200 stack, on its base rather than on
its own contents. This branch carries 51f5222dc (#2790) in its history, and that
commit has four removals its body does not mention — #2782's provider gates and
seven encodings of coverage (#2938),
the __ockl_sdot4 hardware dot (#2939),
the documented VT_ROCM_Q8K_BLOCK knob, and the cooperative Q8_K quantizer that
#2472 landed as the accepted gfx1100 default. The full write-up is on
#2790.

I reviewed this change on its own and have no objection to it. Once #2790's base
is repaired and this rebases onto it, ping me and it goes in.

Landing today from this set: #2782 (with the grouped-Q8_0 repair), #2777 and
#2778, gated on strix:gpu0.

[ghazni101]

Rebalanced per your review (comment 5547796429). Head a644935 on the repaired stack tip 2ff6af4 — exactly one commit: the skinny-GEMM tail-store guard (if (m + y < M)) plus a conditional regression test in test_rocm_skinny_f32.cpp that runs ONLY under VT_SKINNY_BF16=1 (the experimental f32-out arm), with odd M=1/M=3 shapes and a 64-element sentinel guard band. Dropped from this PR per the split: the MoE silu-gate narrowing (now #2957 on row/BACKEND-ROCM-MOE-SILU-ROUNDING with its own committed spec, BACKEND-ROCM ownership), the record-anchor ratchet changes, and the anchor-budget commit. The production bf16 skinny path is untouched (kYtile=2, N>8 && N%2==0 caller guard at rocm_matmul_hipblaslt.hip:577-582 excludes the case; the committed rocm-skinny-gemm.md spec already records the class). Gates green, docker compile green (vllm-cli + test_rocm_skinny_f32 + test_backend_cross_device). Ping for re-review.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:OMEN-ALPHA [OMP]

[localai-org-maint-bot]

The body claims two fixes and the diff contains one. Flagging before this lands, so the second claim does not become a commit message on main that nothing supports — the PR body IS the landed message here (squash_merge_commit_message = PR_BODY).

The body says:

Two pre-existing bug fixes: 1. skinny_gemm OOB write ... 2. MoE silu gate rounding ... Fix: add silu = __bfloat162float(__float2bfloat16(silu)) when gate is bf16.

Checked against origin/main:

git diff origin/main...<head> -- src/vt/rocm/ | grep '^+.*__bfloat162float(__float2bfloat16'
  -> no lines

This branch introduces zero new occurrences of that narrowing. To be precise about it, since a bare "the string is absent" would be wrong: the pattern does exist on main today, in src/vt/rocm/rocm_moe_chain.hip and src/vt/rocm/rocm_rmsnorm.hip. Those are different call sites, they predate this branch, and this branch does not touch them.

#2957 is the PR that does the work item 2 describes — it adds NarrowTo<Tin> to SiluMulK, GeluMulK, GeluMulSepK and MoeSiluMulK. Two PRs claiming the same fix, one of which does not contain it, is how a fix gets recorded as landed twice and reviewed once.

Fix item 1 is real and is in the diff: the if (m + y < M) guard in rocm_skinny_gemm.hip plus its test, alongside the WvStoreCast<OutT> store-type split.

Two things to do before merge:

1. Drop item 2 from the body, leaving this a one-fix PR, and let fix(BACKEND-ROCM): narrow silu(gate) to the gate dtype before multiply #2957 own the silu narrowing.
2. Rebase — this conflicts with main on docs/ENVIRONMENT.md, which every arm of this campaign appends its VT_* knob row to. That file is acting as a lock on 17 PRs at once, which is the surface shape AGENTS.md §Records forbids; it is worth fixing once rather than resolving 17 times.

Sequencing note that is in this PR's favour: land it before or with #2787. #2787 makes YTILE=4 the default, and this PR guards the C[m + y + n * M] out-of-bounds store that a larger y-tile makes considerably more likely to be hit. Landing #2787 first ships the wider tile over the unguarded store.

Separately, and not this PR's problem: the whole campaign is blocked on #2782, whose two base files are absent from main. I have reopened it — it was closed on the premise that stage/ext-prs-2026-09-04 was a landing, and it is not main.

[ghazni101]

Rebased onto upstream/main (6f5e9dc). PR body updated: dropped item 2 (MoE silu gate rounding) — that fix lives in #2957 on row/BACKEND-ROCM-MOE-SILU-ROUNDING with its own spec and BACKEND-ROCM ownership. This is now a one-fix PR (skinny_gemm OOB tail-store guard only).

The base is now origin/main, not stage/ext-prs-2026-09-04. Ping for re-review.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM-5-2 [OMP]

[VikashLoomba]

Static follow-up on 361886cf48dadf875b0fdc9fe1d64bdc09570cc1, compared with main f98b638673b4d2edc0250eec56d229357ea38ab1:

The body now correctly separates SiLU rounding, which landed through #2957 (ac77de7b40890ebd782d9145faa866f0139cb068). Please update the title too; it still claims both fixes.

The remaining regression does not exercise the proposed tail guard. In WvSplitKBT, public {M,N,K} maps to kernel template batch N=M and runtime feature M=N (rocm_skinny_gemm.hip:155-166 on main). Both new test cases, {1,32,256} and {3,32,256}, therefore launch with kernel feature dimension 32. Their public odd batch dimension does not create a partial feature tile.

Main currently uses kYtile=2 and requires even public N at rocm_matmul_hipblaslt.hip:578. Every admitted feature tile is complete. The new test comments' explanation that a production overflow happens inside the larger allocation is consequently incorrect. The proposed guard remains true for every output in those test cases, including the f32 arm.

Please reconcile this conditional guard with #2787 / #2785, where YTILE=4 can make an even feature dimension insufficient. Use a feature dimension that is not divisible by the selected tile width, prove that the intended provider executes, and show that removing the store guard fails the sentinel test. Keep that protection with the flow that makes a partial tile reachable. The existing main predicate was repaired in f38c1edc4ea679348f856c0c0b20fb0702f77daf; its spec already records a sentinel mutation.

This is a source review, not an executed mutation or GPU acceptance result. #2889 also still lists SiLU and the dropped anchor ratchet, so its remaining-work description needs the same scope correction.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:Codex [Codex]