[libratechw]

Title

Issue: none (follows #1372, which added STRATA_GDN_CHUNKED=1 for CUDA; related: #1296, the gfx1100 WY arm)

Summary

STRATA_GDN_CHUNKED=1 exists for CUDA (#1372) and is compiled out on HIP. On gfx1151 the prompt's recurrence is
gdn_rec_quad2c_kernel plus a norm kernel: token by token, FP32, on the vector ALUs. This PR makes the same switch take a gfx11
path: the WY form of the gated delta rule in chunks of 16 tokens, every product one 16x16x16 WMMA (FP16 operands, FP32
accumulation, the state in the waves' accumulators), with the output norm and the z gate fused into the scan.

It is off by default; the default path is unchanged. It rounds differently from the FP32 recurrence (FP16 operands), which is why
it is opt-in; the KL numbers are below.

What changed

• src/prefill/kernels.cu, HIP only, behind the STRATA_GDN_CHUNKED switch in gdn_recurrence_variant (variants 0 and 1, the
  ones gdn_recurrence uses; STRATA_GDN_REC_HEADS is read once at the top of the function now, one declaration moved up):
  • gdn_wmma_prep_kernel: T = (I + A)^-1 and P of every chunk at once (they do not depend on the state); a block per
    (key head, chunk), a wave per value head.
  • gdn_wmma_scan_kernel: a block per value head, a wave per 16 value columns; walks the chunks and loads the next chunk's k/q/v
    while this one computes. Its barriers wait for LDS only (__syncthreads() also waits for every global load in flight, which
    would put the prefetch in front of every barrier), and the prefetched rows are ext_vector_type values (HIP's float4 copies
    go through a stack object, which also waits for the load); both are in the code comments.
  • gdn_rec_wmma: the launch. It returns false, and the existing kernels run, for a part outside gfx_arch_is_gfx11_wmma (gfx1100-1103, gfx1150, gfx1151), fewer
    than 64 tokens and an output row stride not divisible by 8; the caller skips it under STRATA_GDN_REC_HEADS. The part is
    judged per call from the current device of the calling thread. The gfx11-only instructions sit inside the gfx11 device passes.
• tests/hip/gdn_chunked.cpp and CMake: hip_gdn_chunked, hip_gdn_chunked_off, hip_gdn_chunked_default, and
  hip_gdn_chunked --bench / --bench-default (one layer's time on this path / on gfx1151's default kernels). tests/hip/gdn_rec_head.cpp unsets STRATA_GDN_CHUNKED (its variants 0 and 1 are the references there).
• include/strata/prefill/kernels.hpp (the variant comment), docs/AMD_HIP.md (conditions, numbers, how to run the checks),
  docs/DETAILS.md (the CUDA entry now says what HIP does).

Not touched: the CUDA kernels (the #if !defined(__HIPCC__) block), the SYCL port, any default. On HIP STRATA_GDN_CHUNKED=2 is
the same as 1 (the docs say so).

Checks

Ryzen AI Max+ 395 (gfx1151), ROCm 7.14.1, Linux 7.0, UD-IQ4_XS unless stated, all on this commit. "Section 5" is the nine variables of
docs/STRIX_HALO.md section 5 (with STRATA_PF_SWITCH_MIN_T=2048 where the doc has 4096); "defaults" is the engine without them
(section 4's gfx1151 defaults stay).

The default path is unchanged. With the switch unset or 0, for a chunk of fewer than 64 tokens and for a row stride the path
does not take, the output and the carried state are bit-identical to the existing recurrence (hip_gdn_chunked_off,
hip_gdn_chunked_default; both set STRATA_GDN_HEAD=1 and STRATA_GDN_PP=2, gfx1151's defaults).

Numerics (hip_gdn_chunked). 0-16,385 tokens, the 64-token edge, 16-token remainders, row strides 6144 / 6152 / 6145, weak
decay, beta = 0, and a request split at 129 + 128 tokens with the state carried. Bounds: relative RMS of the FP16 output and of the
final state <= 0.003; max |diff| <= 2% of the reference's max + 1e-5; all finite; guard values around the output unchanged. Two
short cases are also compared with a double-precision CPU recurrence of three value heads.

First-token KL against the same run without the switch (STRATA_DUMP_FIRST_LOGITS; 12 natural-text prompts of 1.5K-24K tokens,
nats, full vocabulary):

	UD-IQ4_XS mean / max	UD-Q4_K_XL mean / max	top-1
defaults, with vs without the switch	0.0066 / 0.016	0.0059 / 0.043	12 of 12
section 5, with vs without the switch	0.0043 / 0.013	0.0095 / 0.054	12 of 12

Section 5's own switches against the defaults are 0.0076 (IQ4_XS) and 0.0065 (Q4_K_XL); that is a different change, given for scale
and not as a limit for this one. #1372 reports 0.002-0.009 for the CUDA path. The mean is carried by a few prompts (one at 0.054;
the medians are 0.001-0.004). Only the first generated token is compared, and across the 12 prompts it takes only two distinct values,
so top-1 agreement is a weak check: a long generation, coding / agent use and perplexity have not been measured, which is the
reason for opt-in.

Speed, one layer (hip_gdn_chunked --bench against --bench-default, which runs gfx1151's default kernels with
STRATA_GDN_PP=2 and STRATA_GDN_NOY=1; the switch is read once per process, so the two are separate runs, alternated three
times; each run reports the GPU-event median of 7 runs after 2 warmups with the same nonzero start state; recurrence + norm;
medians of the three):

tokens	default kernels	this PR
64	0.090 ms	0.050 ms	1.8x
256	0.280 ms	0.192 ms	1.46x
2,048	2.603 ms	1.779 ms	1.46x
16,384	20.830 ms	13.203 ms	1.58x
32,768	42.229 ms	26.429 ms	1.60x

Speed, whole prompt. The same binary on both sides, only STRATA_GDN_CHUNKED differing, a fresh engine per side, resident
budget 55 GiB, int8 KV, 11 pairs in alternating order after a discarded one, 256-token answers, the engine's prompt_ms. Every
prompt has its own nonce at the front and was read in full (0 reused tokens in every request). Each engine reads a discarded 4K
prompt first: on those first prompts the same pairs gave 95% intervals of [1.015, 1.021] (section 5) and [0.991, 1.210]
(defaults). Before each prompt the client waits until the page-migration and direct-reclaim counters of /proc/vmstat have
stood still for 3 s; a pair in which either prompt still met page migration (over 50,000 pages) or direct reclaim was dropped
and measured again (2 of 68 pairs, both on the defaults: a 32K prompt with the switch on and a 128K one with it off, which
met 365,000 and 161,000 migrated pages and were 0.51 s and 1.15 s slower than the median of their side; with them kept, 32K
and 128K give 1.011 (10 of 12 and 12 of 12 faster)). Times are
medians; off / on is the geometric mean of the 11 per-pair ratios with a 95% bootstrap interval:

tokens	arm	off -> on	off / on	faster	sign test p
4,096	section 5	3.701 -> 3.630 s	1.019 [1.017, 1.021]	11 of 11	0.001
4,096	defaults	5.947 -> 5.876 s	1.012 [1.008, 1.016]	10 of 11	0.012
32,768	section 5	26.024 -> 25.401 s	1.024 [1.023, 1.026]	11 of 11	0.001
32,768	defaults	44.684 -> 44.122 s	1.012 [1.009, 1.016]	10 of 11	0.012
128,000	section 5	107.468 -> 104.858 s	1.025 [1.023, 1.026]	11 of 11	0.001
128,000	defaults	195.498 -> 193.445 s	1.011 [1.008, 1.013]	11 of 11	0.001

So about 2.4% on top of section 5 and about 1.1% on the defaults, for a prompt that is not the first one after a start. With 36
gated-DeltaNet layers the layer figures predict about 0.55 s saved on a 32K prompt; 0.62 s (section 5) and 0.56 s (defaults) were
measured.

Builds. HIP: kernels.cu compiles for gfx1030, gfx1100, gfx1103, gfx1151, gfx1152 and gfx1201 (the scan kernel holds 26 WMMA
instructions for gfx1100 / gfx1103 / gfx1151 and only a trap for gfx1152 / gfx1201, which the runtime gate never launches it on); the four HIP tests above (the three
hip_gdn_chunked* and hip_gdn_rec_head) pass on gfx1151. Run on gfx1151 only: nothing here has been run on another gfx11 card (gfx1103 joined gfx_arch_is_gfx11_wmma upstream after this
was first written, and this path follows that list) or
with more than one GPU. CUDA (RTX 3060, CUDA 13.0, sm_86, Windows, this commit): strata builds (155 of 155 steps, no error). gdn_rec_parity passes
too, but it is a self-contained copy of the CUDA recurrence that does not link the engine, so it says nothing about this change.
Not built: SYCL (its own sources, none touched).

Extra Notes

• hip: chunk-parallel GDN recurrence (WY form), two parity-gated arms (STRATA_GDN_WY) #1296 adds a WY arm (STRATA_GDN_WY) to the same dispatch for gfx1100, where it loses to the serial kernel. This PR starts from
  gfx1151's gdn_rec_quad2c_kernel, uses 16x16x16 WMMA tiles for every product and reuses the existing switch instead of adding a
  second one. They share no code that I found; whichever lands second needs a small rebase of the dispatch.
• Open question for review: the 64-token minimum and the 16-token chunk come from gfx1151 only. Whether gfx1100 / gfx1150 want the
  same values is untested.

🤖 Generated with Claude Code

[Niko1221]

Ported (two comment/CMake conflicts resolved, your authorship kept as co-author), opt-in only. We cannot run gfx11 matrix cores here; on an R9700 the switch is inert (the new tests report exact equality with the old recurrence, the identity gate is unchanged), and the file builds for gfx1100/1103/1150/1151/1201. Your numbers (+1.4-2.1% for a whole prompt, KL 0.004-0.007) are modest and change bits, so it stays opt-in.

[libratechw]

Thanks for porting this as opt-in. Since you can't run the gfx11 matrix cores, I ran the GDN tests on the merged build (477ac13, gfx1151): ctest -R gdn passes all 5 tests, including hip_gdn_chunked_default and gdn_parity. Nothing further from my side; this PR can be closed.

The test log: https://github.com/libratechw/Strata/blob/09df26f072914f96cfa4fcae3e835e462358808e/main477/gdn-ctest.log

[libratechw]

Thanks for porting this and for the review. I confirmed the change is in main, and there is nothing further I need from this PR, so I'm closing it.