[bong-water-water-bong]

A 2-hour KernelForge campaign (tools/kernelforge/gdn-prefill, about $33 of
Claude usage) sped up the delta-net prefill kernel. ROCmFPX#5 carries the
safe part: DPP reductions, two state columns per lane, bank-interleaved rows
and 16-byte staging. One Qwen3.8-27B layer, 512 tokens: 3.02 -> 0.41 ms
(5.0x with 8 rollback snapshots).

tools/bench.py, Qwen3.8-27B-H32 through 1bit serve, 2 rounds, box lock:
prompt 494 -> 550 tok/s median without a drafter, 467 -> 516 with DFlash2;
decode 42.1 / 24.9 / 13.6 tok/s (code / prose / short) with DFlash2.

The campaign also kept a gate-factored state that turns NaN with hard gates
(log-gates in [-20, 0): ~1.1M non-finite values). My task driver drew gates
from [-0.5, 0), which let it pass; the driver now uses ggml's own test range
[-20, -1e-4], program.md warns against gate products, and the task kernel is
the shipped layout. tools/kernelforge/README.md records the result and two
setup traps (KernelForge's claude --bare ignores a subscription login;
pgrep -f matching its own ssh command). docs/lean.md, serve.md, README and
the site description carry the new numbers.

🤖 Generated with Claude Code

[context7]

Docs7 for 1bit-monster/engine

Result	Status	Action
Deployment	➖ Not used	—
Content review	➖ Did not run. This site has no agent runs available this month. Wait for the monthly reset or check your Docs7 plan.	—

Commit 7f9f116

[github-actions]

PR Reviewer Guide 🔍

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶 5 - Partially compliant Compliant requirements: Remove the CPU reference and tokenizer Port the existing engine instead Remove engine/core, engine/backends/cpu, the tests and golden tools, and their docs Rewrite the README, the porting map and contributing rule 1 Keep the license, NOTICE, the copyright check and CI Non-compliant requirements: None

⏱️ Estimated effort to review: 3 🔵🔵🔵⚪⚪

🧪 No relevant tests

🔒 No security concerns identified

⚡ Recommended focus areas for review Gate Range Change The gate range has been changed from [-0.5, 0) to [-20, -1e-4) to better align with ggml's test range. This change is important to prevent underflow issues with hard gates that could lead to NaN values in the state, which was previously masked by a gentler range in the driver. # log decay over ggml's own test range (test-backend-ops GATED_DELTA_NET): hard gates make any # product of gates underflow, which a gentler range let a campaign miss (tools/kernelforge/README.md) g = -1e-4 - torch.rand(T, H, generator=gen) * (20.0 - 1e-4) Kernel Optimization The kernel has been optimized with DPP reductions, two state columns per lane, bank-interleaved rows, and 16-byte staging. These changes significantly improve performance, reducing the prefill kernel time from 3.02 ms to 0.41 ms for a 512-token layer. // XOR-reduce across the 4 lanes of a quad and broadcast. __shfl_xor lowers to ds_bpermute on RDNA, // which issues on the *LDS pipe* -- the same pipe the k_s/q_s reads contend for, and it needs an // lgkmcnt drain that also retires those reads, dragging them onto the token recurrence. DPP // quad_perm is a VALU modifier instead: no LDS traffic and no wait. template <int WIDTH> __device__ __forceinline__ float group_sum(float x) { static_assert(WIDTH == 4 || WIDTH == 8, "DPP reduction covers SPLIT 4 and 8"); // row_half_mirror is the XOR-7 butterfly over each group of 8 lanes; composing it with the two // quad_perms covers {1,2,3,4,5,6,7}, i.e. a full 8-lane reduce-and-broadcast in three DPP adds. if constexpr (WIDTH == 8) x += __builtin_bit_cast(float, __builtin_amdgcn_update_dpp( 0, __builtin_bit_cast(int, x), 0x141, 0xF, 0xF, true)); // quad_perm[1,0,3,2] = XOR 1, quad_perm[2,3,0,1] = XOR 2. x += __builtin_bit_cast(float, __builtin_amdgcn_update_dpp( 0, __builtin_bit_cast(int, x), 0xB1, 0xF, 0xF, true)); x += __builtin_bit_cast(float, __builtin_amdgcn_update_dpp( 0, __builtin_bit_cast(int, x), 0x4E, 0xF, 0xF, true)); return x; } // Lane -> state-row map. The blocked map (row = part * ROWS) puts the SPLIT lanes of a column on // LDS dwords 0, 32, 64, 96 of the same token row, i.e. all four on bank 0: a 4-way conflict on // every k_s/q_s read. Interleaving in groups of 4 instead -- // row(r) = (r >> 2) * (SPLIT * 4) + part * 4 + (r & 3) // -- is still a bijection onto 0..S_v-1, still gives each lane four *consecutive* dwords (so the // ds_load_b128 survives, base part*16, immediates 64 * (r >> 2)), but spreads the four lanes over // banks 4 * part .. 4 * part + 3, which is conflict-free. Row padding cannot fix the blocked map: // part * ROWS is a multiple of 32 for any row stride. The state tensor layout is unchanged -- // s[r] simply denotes a different row -- so only the partial-sum partition of the reduction moves. #define GDN_ROW(r) ((((r) >> 2) * (SPLIT * 4)) + part * 4 + ((r) & 3)) Gate Underflow Prevention The program documentation now explicitly warns against factoring gates out of the state due to underflow issues with hard gates, which could lead to NaN values. This is a critical safety measure to ensure numerical stability. - Gates are hard: log-gates reach -20 (the driver draws them from [-20, -1e-4], as ggml's test does). A product of gates over a few tokens underflows float32, so do not factor the gates out of the state (S = prod(g) * S_hat): an earlier campaign's kept candidate did and returned NaN. - The current kernel already has the first campaign's layout (DPP reductions, 2 columns per lane, interleaved rows, 16-byte staging); its K=1 case runs in about 0.41 ms.