[bong-water-water-bong]

A second KernelForge campaign (tools/kernelforge/gemv-q4_0, about $40) on the
Q4_0 x Q8_1 few-column matmul that speculative verification runs found a
row-tiled body. ROCmFPX#6 dispatches it for 8-column batches on gfx1151: its
output is bit for bit the old kernel's (40 shape/input cases), and Qwen3.8-27B
8-token batches go 71.3 -> 78.3 tok/s.

tools/bench.py, Qwen3.8-27B-H32 through 1bit serve, 2 rounds, box lock:
base 559 / 546 tok/s prompt, 13.1 / 13.1 / 13.5 decode; --dflash 518 / 486
prompt, decode code / prose / short 46.0 / 28.4 / 17.6 (best; medians 45.6 /
28.0 / 17.5), with the rocm p-min 0.4 recipe applied. Before these two changes
the same tool measured 42.1 / 24.9 / 13.6.

tools/kernelforge/gemv-q4_0: the task (kernel now at the kept body, float64
oracle with adversarial inputs, driver, program.md with the first campaign's
findings, run.sh); README records the result. docs/lean.md, serve.md, README
and the site description carry the new numbers; serve.md's p-min sentence
follows the recipes.

🤖 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 5ebb976

[github-actions]

Preparing review...

[github-actions]

Preparing review...

[github-actions]

Preparing review...

[github-actions]

PR Reviewer Guide 🔍

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶 6 - Partially compliant Compliant requirements: Embedded Lemonade with v11.9.0, complete, plus local deltas in its UPSTREAM.md 1bit lemonade runs Lemonade's server in-process Smoke test (CI): health, catalog and system info Pinned sources: ROCm/hrx 0bc22fb, ROCm/hrx-system 6743075f, and llama.cpp fork on branch 1bit-engine/hrx-vulkan Build with -DONEBIT_HRX=ON builds one llama-server with HRX20 and Vulkan0 Verified on Strix Halo with 48/48 kernel artifacts byte-identical to working build's Decode speed on par with working build Lemonade serves same checkpoint on Vulkan0 (117 tok/s) and HRX20 (34 tok/s), and both answer correctly CI covers default build (ONEBIT_HRX=OFF) and smoke test Non-compliant requirements: HRX runs on Strix Halo only, since GitHub runner has no AMD GPU Requires further human verification: HRX runs on Strix Halo only, since GitHub runner has no AMD GPU

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

🧪 No relevant tests

🔒 No security concerns identified

⚡ Recommended focus areas for review Possible Performance Regression The driver.py file introduces adversarial input generation for testing correctness, but it may inadvertently cause performance issues or unexpected behavior during benchmarking if not carefully managed. Specifically, the use of rng.random(q.shape) in the quantize_q4_0 function with mode "edges" could lead to inconsistent performance characteristics across different runs, which might mask or amplify performance regressions in the kernel. q = np.where(rng.random(q.shape) < 0.5, 0, 15).astype(np.uint8) packed = (q[..., :16] | (q[..., 16:] << 4)).astype(np.uint8) Kernel Code Complexity The kernel implementation in gemv_q4_0_kernel.py is highly complex with multiple template specializations and conditional logic for handling different column counts. While this is necessary for performance optimization, it increases the risk of subtle bugs or incorrect behavior in edge cases, especially when dealing with different memory access patterns and register usage. template <int NCOLS, int R> __global__ void __launch_bounds__(WARP, 1) gemv_q4_0_rt(const block_q4_0 * __restrict__ x, const block_q8_1 * __restrict__ y, float * __restrict__ dst, int ncols_x, int nrows) { static_assert(R * NCOLS <= WARP, "one lane per result in the epilogue"); const int tid = threadIdx.x; const int blocks_per_row = ncols_x / 32; constexpr int blocks_per_iter = VDR * WARP / QI4_0; // 16 const int row0 = blockIdx.x * R; // consecutive rows: one narrow weight window const int kqs = VDR * (tid % (QI4_0 / VDR)); float tmp[R][NCOLS] = {{0.0f}}; const block_q4_0 * xr = x + (size_t) row0 * blocks_per_row; for (int kbx = tid / (QI4_0 / VDR); kbx < blocks_per_row; kbx += blocks_per_iter) { int u[NCOLS][2 * VDR]; // this lane's slice of every column's k-block: read once float2 ds[NCOLS]; #pragma unroll for (int j = 0; j < NCOLS; ++j) { const block_q8_1 * bq8 = &y[(size_t) j * blocks_per_row + kbx]; #pragma unroll for (int i = 0; i < VDR; ++i) { u[j][2 * i + 0] = get_int_b4(bq8->qs, kqs + i); u[j][2 * i + 1] = get_int_b4(bq8->qs, kqs + i + QI4_0); } ds[j] = __half22float2(bq8->ds); } #pragma unroll for (int r = 0; r < R; ++r) { const block_q4_0 * bq4 = &xr[(size_t) r * blocks_per_row + kbx]; int vi0[VDR], vi1[VDR]; #pragma unroll for (int i = 0; i < VDR; ++i) { // unpacked once, used by all NCOLS columns const int vv = get_int_b2(bq4->qs, kqs + i); vi0[i] = (vv >> 0) & 0x0F0F0F0F; vi1[i] = (vv >> 4) & 0x0F0F0F0F; } const float dw = __half2float(bq4->d); #pragma unroll for (int j = 0; j < NCOLS; ++j) { int sumi = 0; #pragma unroll for (int i = 0; i < VDR; ++i) { sumi = dp4a(vi0[i], u[j][2 * i + 0], sumi); sumi = dp4a(vi1[i], u[j][2 * i + 1], sumi); } tmp[r][j] += dw * (sumi * ds[j].x - (8 * VDR / QI4_0) * ds[j].y); } } } // Epilogue. Every lane holds a partial for each of the R*NCOLS results and the warp needs, per // result, the sum over lanes. When there is exactly one result per lane, a reduce-scatter // halves the live accumulator count at each xor stage instead of running an independent // 5-stage butterfly per accumulator: 31 shuffles for 32 results rather than 160, and lane i // ends holding result i, which also deletes the 32-way select chain before the store. The // pairing order (xor 16, 8, 4, 2, 1) is the butterfly's, so each result is summed by the same // balanced tree over the same partials -- bit for bit the same float. if constexpr (R * NCOLS == WARP) { float acc[WARP]; #pragma unroll for (int j = 0; j < NCOLS; ++j) { #pragma unroll for (int r = 0; r < R; ++r) acc[j * R + r] = tmp[r][j]; } #pragma unroll for (int m = WARP / 2; m > 0; m >>= 1) { const bool hi = (tid & m) != 0; #pragma unroll for (int k = 0; k < m; ++k) { const float a = acc[k], b = acc[k + m]; acc[k] = (hi ? b : a) + __shfl_xor(hi ? a : b, m, WARP); } } dst[(size_t) (tid / R) * nrows + row0 + (tid % R)] = acc[0]; } else { #pragma unroll for (int j = 0; j < NCOLS; ++j) { #pragma unroll for (int r = 0; r < R; ++r) { #pragma unroll for (int o = WARP / 2; o > 0; o >>= 1) tmp[r][j] += __shfl_xor(tmp[r][j], o, WARP); } } if (tid < R * NCOLS) { // lane j*R + r holds column j of row row0 + r float v = tmp[0][0]; #pragma unroll for (int j = 0; j < NCOLS; ++j) { #pragma unroll for (int r = 0; r < R; ++r) if (tid == j * R + r) v = tmp[r][j]; } dst[(size_t) (tid / R) * nrows + row0 + (tid % R)] = v; } } } // x: nrows * (ncols_x/32) block_q4_0; y: ncols * (ncols_x/32) block_q8_1; dst: [ncols][nrows] float. extern "C" void gemv_q4_0_hip(const void * x, const void * y, float * dst, int ncols_x, int nrows, int ncols, hipStream_t stream) { const dim3 grid(nrows), block(WARP); const block_q4_0 * bx = (const block_q4_0 *) x; const block_q8_1 * by = (const block_q8_1 *) y; // Row tiling only pays where the activation side is the scaling cost; 1..4 columns already run // at ~88% of LPDDR5X peak, so they keep the one-row body verbatim (no risk to plain decode). constexpr int R = ROWS_PER_WARP; if (ncols >= 5 && R > 1 && nrows % R == 0) { const dim3 grid_rt(nrows / R); switch (ncols) { case 5: hipLaunchKernelGGL((gemv_q4_0_rt<5, R>), grid_rt, block, 0, stream, bx, by, dst, ncols_x, nrows); return; case 6: hipLaunchKernelGGL((gemv_q4_0_rt<6, R>), grid_rt, block, 0, stream, bx, by, dst, ncols_x, nrows); return; case 7: hipLaunchKernelGGL((gemv_q4_0_rt<7, R>), grid_rt, block, 0, stream, bx, by, dst, ncols_x, nrows); return; case 8: hipLaunchKernelGGL((gemv_q4_0_rt<8, R>), grid_rt, block, 0, stream, bx, by, dst, ncols_x, nrows); return; default: break; } } switch (ncols) { case 1: hipLaunchKernelGGL(gemv_q4_0<1>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 2: hipLaunchKernelGGL(gemv_q4_0<2>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 3: hipLaunchKernelGGL(gemv_q4_0<3>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 4: hipLaunchKernelGGL(gemv_q4_0<4>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 5: hipLaunchKernelGGL(gemv_q4_0<5>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 6: hipLaunchKernelGGL(gemv_q4_0<6>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 7: hipLaunchKernelGGL(gemv_q4_0<7>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; case 8: hipLaunchKernelGGL(gemv_q4_0<8>, grid, block, 0, stream, bx, by, dst, ncols_x, nrows); break; default: break; } } Execution Environment Assumptions The run.sh script assumes that the Hyperloom environment is properly configured with the kernelforge CLI and claude CLI available on the PATH. If these dependencies are not correctly installed or configured, the script will fail. This could be problematic in CI environments or when running on different systems. command -v claude >/dev/null || { echo "run.sh: no claude CLI on PATH; KernelForge drives its agent through it"; exit 1; }