[bsorensen110]

Problem and result

Supersedes #3064 (closed during the 0.2.0-beta tracker reset), rebased onto 0.2.0-beta (e88ee39d).

In src/kernels/cuda/verify_kernels.cu (the GDN speculative-verification and decode kernels launched 36 times per verify window):

1. gdn_conv_l2_multi_kernel and gdn_step_norm_groups_kernel used full-workgroup __syncthreads() (s_wait_loadcnt_dscnt 0x0; s_barrier; global_inv scope:SCOPE_SE on gfx12) where only LDS visibility (s_cs, s_w, s_qkv, s_k, s_v, s_d, s_out) is required, draining in-flight global loads and invalidating the L0 vector cache inside the recurrence loop.
2. When STRATA_GDN_SPLIT=1 is enabled, gdn_conv_l2_multi with commit_state=false left conv_state commitment to Verifier::commit, which previously launched 36 separate single-layer gdn_conv_commit kernels per verify window (!gdn_split_on() gated off the 1-launch batched gdn_conv_commit_batch_kernel path in verify.cpp unless STRATA_GDN_COMMIT_SPLIT was explicitly set).

This PR adds LDS-only vk_sync_lds() for gdn_conv_l2_multi_kernel and gdn_step_norm_groups_kernel, enables gdn_commit_split_on() automatically whenever gdn_split_on() is active (STRATA_GDN_COMMIT_SPLIT=0 overrides), and extends src/kernels/gdn_parity.cpp with multi-token gdn_conv_l2_multi + gdn_step_norm_multi (groups=0,1,2,4) + gdn_conv_commit_batch bitwise parity tests across T = 1..8. All outputs are 100% bitwise identical (0 differing bits).

Related discussion

Supersedes #3064. Default dispatch when STRATA_GDN_SPLIT is unset is unchanged (vk_sync_lds() optimizes the default gdn_conv_l2_multi_kernel and gdn_step_norm_groups_kernel paths without altering any bits).

Approach

• Define vk_sync_lds() (s_wait_dscnt 0x0; s_barrier_signal -1; s_barrier_wait -1 on gfx12, s_waitcnt lgkmcnt(0); s_barrier on gfx1030/gfx11, __syncthreads() fallback) in src/kernels/cuda/verify_kernels.cu.
• Update gdn_commit_split_on() to default to gdn_split_on() when STRATA_GDN_COMMIT_SPLIT is unset.
• Add multi-token bitwise parity checks in src/kernels/gdn_parity.cpp covering commit_state=true vs commit_state=false + gdn_conv_commit / gdn_conv_commit_batch and gdn_step_norm_multi (groups = 0, 1, 2, 4) across T = 1..8.

Validation

Passed:

• cmake --build build-copilot-hip --target strata gdn_parity arch_defaults_test -j16: revision 1284c5e4 on 0.2.0-beta (e88ee39d), AMD Radeon AI PRO R9700 (gfx1201, 32 GB, Linux 7.2.9-2-cachyos-server, ROCm 10.2.0a20260914), clean build & link.
• ./build-copilot-hip/arch_defaults_test: revision 1284c5e4 (0.2.0-beta), arch_defaults_test: OK.
• ./build-copilot-hip/gdn_parity and STRATA_GDN_SPLIT=1 ./build-copilot-hip/gdn_parity: revision 1284c5e4 (0.2.0-beta), gfx1201, gdn multi-token verify parity (conv_l2_multi, conv_commit_batch, step_norm_groups 1/2/4 vs 0, T=1..8): pass (0 bit diffs), gdn: 0 failures, gdn_parity OK.

Failed:

• none observed

Not run:

• Target-device runtime execution on gfx1151 and CUDA hardware (no gfx1151 or NVIDIA GPU in this host).
• SYCL and Metal builds (verify_kernels.cu is CUDA/HIP only).

Limits

• STRATA_GDN_SPLIT remains off by default on gfx12 (STRATA_GDN_SPLIT=1 is opt-in on gfx12 and default on gfx1151).

---

Summary

Optimizes the Gated Delta Net (GDN) speculative-verify forward and commit kernels in src/kernels/cuda/verify_kernels.cu while preserving 100% bit-for-bit output and state parity:

1. Enable 2-launch batched gdn_commit_layers when STRATA_GDN_SPLIT=1 (verify_kernels.cu:1073):

   • gdn_commit_layers previously checked (sp && sp[0] == '1') (STRATA_GDN_SPLIT=1) and returned false, forcing verify.cpp:2419 to fall back to 72 per-layer kernel launches (gdn_conv_commit + gdn_state_replay_kernel across 36 GDN layers) instead of the 2-launch 36-layer batched commit (gdn_conv_commit_layers_kernel + gdn_step_commit_layers_kernel).
   • Because gdn_step_commit_layers_kernel and gdn_state_replay_kernel share the exact same 32 × 4 column-split state layout and red[0..3][col_local] reduction order, their committed states are 100% bitwise identical (gdn_parity Section 6: 8 n_keep cases × 5 layers, 0 mismatches).
   • Notably, STRATA_GDN_SPLIT=1 is enabled by default on gfx1151 (Strix Halo) in src/core/arch_defaults.cpp, so gfx1151 was silently falling back to 72 per-layer commit launches on every verify window.
   • STRATA_GDN_COMMIT_SPLIT=0 and STRATA_DBG_GDN=1 remain available as explicit per-layer commit overrides, and gdn_parity.cpp now fails if gdn_commit_layers is unexpectedly skipped when neither override is set.

2. LDS-only workgroup barriers (vk_sync_lds()) in GDN verify kernels (verify_kernels.cu:37):

   • On RDNA4 (gfx1200/gfx1201) and RDNA3 (gfx11xx/gfx1030), __syncthreads() compiles to a full workgroup memory fence (s_wait_loadcnt 0x0; s_wait_storecnt 0x0; s_wait_dscnt 0x0; s_barrier_signal -1; s_barrier_wait -1; global_inv scope:SCOPE_SE on gfx12), draining all in-flight global memory loads and invalidating the shader engine's L1 vector cache at every LDS synchronization point (4× per token in gdn_step_norm_multi_kernel, 5× per token in gdn_step_split_kernel, 3× per token in gdn_step_commit_body).
   • Since inter-thread communication in gdn_conv_l2_body, gdn_step_commit_body, gdn_step_norm_multi_kernel, gdn_step_split_kernel, gdn_state_replay_kernel, and gdn_out_norm_kernel goes strictly through __shared__ memory (LDS), vk_sync_lds() uses s_wait_dscnt 0x0; s_barrier_signal -1; s_barrier_wait -1 on gfx12 and s_waitcnt lgkmcnt(0); s_barrier on gfx11/gfx1030 (falling back to __syncthreads() on CUDA and other targets).

3. Prologue state-load ordering and unified next-token prefetch in gdn_step_norm_multi_kernel (verify_kernels.cu:328):

   • Issuing the 32 per-thread state loads (s[0..31]) before the initial sk[0]/sq[0] LDS store allows vk_sync_lds() at t = 0 to overlap all 32 global state loads with the first barrier.
   • Unifying the rg == 0 (sk) and rg == 1 (sq) next-token prefetch under a single predicate (load_kq) overlaps the global load of token t + 1's k/q and token t's v/z with the 32-FMA kv and o loops while reducing register pressure by 4–7 VGPRs across all four gdn_step_norm_multi_kernel instantiations.

4. Double-buffered 3-barrier gdn_step_split_kernel (verify_kernels.cu:417):

   • gdn_step_split_kernel<true> previously used single-buffered sk[S], sq[S], and red[RG][GS_COLS], requiring 5 barriers per token and stalling on ht global loads inside the barrier pair.
   • Double-buffering sk[2][S] and sq[2][S] with separate red_kv[RG][GS_COLS] and red_o[RG][GS_COLS] (LDS 1536 B -> 3072 B, preserving full 16-wave / 4-CTA-per-WGP occupancy) reduces the loop to 3 vk_sync_lds() barriers per token and overlaps next-token k/q global loads with the 32-FMA kv dot product.

---

Root Cause & RED / GREEN Verification

On origin/main (477ac13c), running gdn_parity with STRATA_GDN_SPLIT=1 skipped Section 6 (gdn_commit_layers == per-layer commit) because gdn_commit_layers returned false. With the tightened assertion in src/kernels/gdn_parity.cpp:

• RED (origin/main verify_kernels.cu + STRATA_GDN_SPLIT=1 ./gdn_parity):

  gdn_ab_multi T=1..8 == fused_gdn_ab        bitwise (8 cases)
  gdn_step_norm_multi T=1..8 == single-token bitwise (47 cases)
  gdn_commit_layers == per-layer commit      *** NO (unexpectedly skipped) ***
  
  gdn: 1 failures
  

• GREEN (this PR + STRATA_GDN_SPLIT=1 ./gdn_parity):

  gdn_ab_multi T=1..8 == fused_gdn_ab        bitwise (8 cases)
  gdn_step_norm_multi T=1..8 == single-token bitwise (47 cases)
  gdn_commit_layers == per-layer commit      bitwise (8 cases, 5 layers)
  
  gdn: 0 failures
  

• STRATA_GDN_SPLIT=0 ./gdn_parity, STRATA_GDN_COMMIT_SPLIT=0 ./gdn_parity, ./verify_batch_parity --selftest, ./spec_verify_parity, and ./arch_defaults_test all pass with 0 failures.

---

Code Object Resource Usage (gfx1201, llvm-readobj --notes)

Kernel (gfx1201)	origin/main (VGPR / SGPR / LDS / Spills)	This PR (VGPR / SGPR / LDS / Spills)
gdn_step_norm_multi_kernel<true, false>	89 / 32 / 6160 B / 0	83 / 32 / 6160 B / 0 (-6 VGPRs)
gdn_step_norm_multi_kernel<true, true>	92 / 38 / 6160 B / 0	88 / 38 / 6160 B / 0 (-4 VGPRs)
gdn_step_norm_multi_kernel<false, false>	168 / 38 / 6160 B / 0	162 / 34 / 6160 B / 0 (-6 VGPRs, -4 SGPRs)
gdn_step_norm_multi_kernel<false, true>	173 / 38 / 6160 B / 0	166 / 38 / 6160 B / 0 (-7 VGPRs)
gdn_step_split_kernel<true>	140 / 30 / 1536 B / 0	142 / 30 / 3072 B / 0 (5 -> 3 barriers/tok)
gdn_step_commit_layers_kernel	141 / 28 / 1024 B / 0	141 / 28 / 1024 B / 0 (LDS-only barriers)

Compiled cleanly with zero errors across HIP targets (gfx1100, gfx1101, gfx1102, gfx1103, gfx1151, gfx1200, gfx1201) and CUDA targets (sm_75, sm_89, sm_100, sm_120).

---

Benchmark Evidence (AMD Radeon AI PRO R9700 gfx1201, ROCm 10.2)

1. Interleaved 36-Layer GDN Microbenchmark (5 interleaved pairs)

Geometry matches Qwen3.8-Flash-Next 397B (S=128, h_k=16, h_v=48, C=10240, value_dim=6144, MT=8, L=36 layers; 117.7 MB total GDN state across 36 layers so every window streams cold from VRAM beyond the 64 MB Infinity Cache). Each trial averages 50 full 36-layer verify windows after 10 warmup windows.

STRATA_GDN_SPLIT=0 (default)

Operation (36 layers total, us)	origin/main median (range)	This PR median (range)	Median Delta
forward_unfused (T=1)	452.32 (451.92..453.15)	445.76 (445.07..445.82)	-6.56 us (-1.5%)
forward_unfused (T=2)	505.54 (504.86..506.28)	491.39 (490.32..491.66)	-14.15 us (-2.8%)
forward_unfused (T=4)	613.94 (613.03..614.64)	586.45 (585.36..587.67)	-27.49 us (-4.5%)
forward_unfused (T=5)	667.81 (667.07..668.67)	634.13 (633.48..635.22)	-33.68 us (-5.0%)
forward_qfuse (T=1)	342.04 (341.81..342.26)	337.62 (337.37..337.64)	-4.42 us (-1.3%)
forward_qfuse (T=2)	401.96 (401.82..402.47)	389.40 (387.85..390.58)	-12.56 us (-3.1%)
forward_qfuse (T=4)	523.24 (522.12..523.72)	497.17 (495.51..498.18)	-26.07 us (-5.0%)
forward_qfuse (T=5)	580.32 (578.97..580.56)	557.36 (555.50..559.85)	-22.96 us (-4.0%)
commit_per_layer (nk=5)	768.05 (768.05..768.84)	759.76 (758.62..759.97)	-8.29 us (-1.1%)
commit_batched (nk=5)	422.14 (421.90..422.33)	420.01 (419.88..420.10)	-2.13 us (-0.5%)

STRATA_GDN_SPLIT=1 (also gfx1151 default in arch_defaults.cpp)

Operation (36 layers total, us)	origin/main median (range)	This PR median (range)	Median Delta
forward_unfused (T=1)	452.37 (452.02..452.97)	446.02 (444.74..447.13)	-6.35 us (-1.4%)
forward_unfused (T=2)	593.98 (593.03..594.65)	581.03 (580.32..581.24)	-12.95 us (-2.2%)
forward_unfused (T=4)	695.61 (694.78..696.42)	657.51 (656.41..658.53)	-38.10 us (-5.5%)
forward_unfused (T=5)	746.94 (746.28..748.58)	698.87 (697.13..700.95)	-48.07 us (-6.4%)
forward_qfuse (T=1)	341.94 (341.64..342.27)	337.60 (337.59..337.75)	-4.34 us (-1.3%)
forward_qfuse (T=2)	483.98 (482.91..484.31)	470.88 (470.36..471.35)	-13.10 us (-2.7%)
forward_qfuse (T=4)	584.12 (583.46..585.01)	548.27 (546.97..548.80)	-35.85 us (-6.1%)
forward_qfuse (T=5)	641.35 (639.26..642.32)	591.03 (589.98..591.54)	-50.32 us (-7.8%)
Effective 36L commit (nk=1)	677.49 (674.59..678.83, 72-launch fallback)	456.33 (456.09..456.87, 2-launch batched)	-221.16 us (-32.6%, 1.48x)
Effective 36L commit (nk=2)	710.57 (710.15..711.48, 72-launch fallback)	451.44 (450.61..452.58, 2-launch batched)	-259.13 us (-36.5%, 1.57x)
Effective 36L commit (nk=3)	725.54 (724.37..730.42, 72-launch fallback)	448.58 (443.12..452.33, 2-launch batched)	-276.96 us (-38.2%, 1.62x)
Effective 36L commit (nk=4)	749.85 (748.71..750.47, 72-launch fallback)	415.29 (415.17..418.40, 2-launch batched)	-334.56 us (-44.6%, 1.81x)
Effective 36L commit (nk=5)	771.88 (770.62..778.58, 72-launch fallback)	420.15 (419.96..420.41, 2-launch batched)	-351.73 us (-45.6%, 1.84x)

2. End-to-End 397B --serve Decode A/B (gfx1201)

Workload: Swift-Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S (--spec 4 --spec-min-p 0.5 --mtp .../rt --pcie-frac 0.30 --adapt-every 0 --kv int8 --prefill 4096 --prompt-cache 0, 512 prompt tokens, 256 generated tokens, 1 warmup + 3 measured trials per arm, 15s cooldown between arms). All arms use STRATA_HIPBLASLT_TUNING=tools/hip/gfx1201-hipblaslt-100500.txt STRATA_DENSE_MMQ=1 STRATA_PF_FUSED=1 STRATA_HC_Q8=1 STRATA_HIP_WMMA=1.

Arm	Median Decode (tok/s)	Median Decode (ms)	Per-Trial Decode (tok/s)	Exact Token Match (768/768)
origin/main (STRATA_GDN_SPLIT=0, STRATA_QFUSE=0)	91.33	2803.1	[91.61, 91.33, 91.18]	True (ref)
This PR (STRATA_GDN_SPLIT=0, STRATA_QFUSE=0)	91.97 (+0.7%)	2783.6 (-19.5 ms)	[92.13, 91.97, 91.36]	True
This PR (STRATA_GDN_SPLIT=0, STRATA_QFUSE=1)	92.12 (+0.9%)	2779.0 (-24.1 ms)	[92.23, 92.12, 91.51]	True
This PR (STRATA_GDN_SPLIT=1, STRATA_QFUSE=0)	91.66	2792.8	[91.76, 91.66, 91.53]	True

(Note: Because the 36-layer GDN recurrence step accounts for ~0.65 ms of a ~25.0 ms 397B verify window on R9700 when experts stream over PCIe/CPU, the -34 us to -352 us kernel-level savings translate to a modest ~0.4–0.9% end-to-end wall-time improvement on 397B with overlapping trial ranges on a 3-trial sample; the primary isolated evidence is the 5-pair interleaved 36-layer GDN microbenchmark above.)

---

What This Does NOT Fix / Scope

• Does not change default environment switches in src/core/arch_defaults.cpp (STRATA_GDN_SPLIT and STRATA_QFUSE remain opt-in on gfx12 and default-on on gfx1151).
• Does not modify GDN prompt prefill kernels (src/kernels/cuda/gdn.cu) or MTP / S2 grouped expert kernels.
• On gfx11xx (gfx1151 Strix Halo) and CUDA targets (sm_75, sm_89, sm_100, sm_120), compilation is verified; target-device runtime validation on gfx1151 and CUDA hardware was not run on this gfx1201 host.

Reproduction Commands

cmake --build build-copilot-hip --target gdn_parity verify_batch_parity spec_verify_parity arch_defaults_test -j16
STRATA_GDN_SPLIT=0 ./build-copilot-hip/gdn_parity
STRATA_GDN_SPLIT=1 ./build-copilot-hip/gdn_parity
STRATA_GDN_COMMIT_SPLIT=0 ./build-copilot-hip/gdn_parity
./build-copilot-hip/verify_batch_parity --selftest
./build-copilot-hip/spec_verify_parity
./build-copilot-hip/arch_defaults_test

[Copilot]

🔵 Needs a closer look

Architecture-specific synchronization changes need human review, with runtime validation still missing on gfx1151 and CUDA hardware.

No open issues 🎉

What changed in this PR

Optimizes GDN speculative verification and state commits, aiming to preserve bitwise parity while reducing synchronization and kernel-launch overhead.

Changes:

• Adds LDS-only barriers on selected AMD targets and improves prefetching and buffering.
• Enables batched commits with STRATA_GDN_SPLIT=1.
• Makes parity checks fail when batched commits are unexpectedly skipped.

File	Description
src/​kernels/​gdn_parity.cpp	Checks that batched commits run unless explicitly disabled.
src/​kernels/​cuda/​verify_kernels.cu	Optimizes synchronization, token buffering, and batched commit selection.

_{Review effort: Balanced}