[laudney]

Part of #1977. Stacked on #1978 (the Qwen3 dense HIP core); until #1978 lands, this diff also shows its commits. The reranker numbers need #1981's route.

On RDNA4 (gfx12), Qwen3-8B's BF16 projections can run on a new WMMA tile, opt-in with --policy matrix=on. The default matrix=auto keeps the dot2 tile and its bits.

AMD Radeon AI PRO R9700, cold, median of 5	dot2 (default, today)	WMMA (--policy matrix=on, this PR)	SGLang (same card)
Qwen3-Embedding-8B, 100 short inputs in one request	2.470 s	0.985 s	0.841 s
Qwen3-Embedding-8B, one short query	75.1 ms	40.5 ms	38.4 ms
Qwen3-Reranker-8B, 16 documents of ~260 tokens, one request	3.218 s	1.389 s	0.534 s
One layer's dense products at 16 / 276 / 2,048 rows	1,884 / 4,742 / 26,492 µs	1,006 / 1,912 / 8,554 µs

Accuracy	dot2	WMMA
Layer-0 projections against float64, BF16 out (what the engine writes), q k v o gate up down	1.63e-3 to 1.69e-3	the same: ratio 1.0000 in every class
Layer-0 projections against float64, fp32 sums	~1e-7 relative	8e-7 to 8e-6 (8 to 41 times further)
Embedding oracle (check, #1978's 1× limits): worst distance at width 4096 / 1024 / 32	1.32e-4 / 1.18e-4 / 7.92e-5	1.31e-4 / 1.35e-4 / 1.06e-4 (limits 1.6e-4 / 1.7e-4 / 1.9e-4)
Reranker against transformers float32 (#1981's oracle): worst |yes − no|, pairs swapped	0.084, 2	0.075, 0
Reranker against transformers BF16: worst |score| (gate 0.02)	0.0082	0.0159
Kernel contract on gfx1201, fp32 out: worst error as a share of the bound		1.4% (both geometries, gates 1-4 pass); the planted control (each 16-deep step summed in BF16) red at gate 4

Arithmetic rule (CONTRIBUTING). WMMA changes the sum order, and gfx12's WMMA sums each 16-product step less precisely than the dot2 chain: the fp32 sums are 8 to 41 times further from float64 on Qwen3-8B's real layer-0 operands. The BF16 outputs the engine writes are as close as dot2's, and both models pass their oracles. On review, WMMA is opt-in: only --policy matrix=on serves it, so the default output keeps dot2's bits. gfx1200 builds the same tile but is unmeasured: no run, and no gate 5 record (the records cover gfx1201).

Why

The BF16 projections are nearly all of a Qwen3-8B prompt pass. On the dot2 tile they reach 22 to 30 TFLOPS. The WMMA tile reaches 45 to 67 at 276 rows and 86 to 99 at 2,048 rows. With it, embeddings come within 18% of SGLang. The reranker's remaining gap is one pass per document (about 276 rows each); packing a request's documents into one pass is the next step, separately.

What this changes

• Kernel. zig/kernels/hip/tiles/wmma12_gemm.hpp: y = x · Wᵀ for BF16 x and BF16 W with fp32 accumulators on __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12 (layout checked against the CPU). Two block geometries, a 16 x 64 block (K steps of 64) for few rows or narrow outputs and a 128 x 128 block (K 32, band launch order) above; the launcher picks by size (shapeOf). Both add each output's products in k order, 16 a step, so a row's bytes depend on neither m nor the geometry; the kernel tests check that at every row count around the thresholds. It traps on chips without gfx12 WMMA.
• Selection. The dense kernels take the policy's matrix choice at load: WMMA on gfx12 under on only, dot2 under auto and off and on every other chip. Only qwen3 loads BF16 dense weights, so no other family's bits change. matrix=on is admitted on gfx12 for qwen3 (all its products have the tile); the MLX-affine families keep the refusal.
• Contract. zig/tests/hip/contract_dense_wmma12.zig registers each geometry, fp32 and BF16 out, as a new kernel (not a replaces entry): depth k plus a k 2⁻¹²⁶ floor, the served Qwen3-8B shapes and both blocks' edges; gate-5 records for gfx1201 from ROCm 10.1. Its planted control, tf_ctl_wmma12_bf16_sum in the controls' code object, is the small block with each step through gfx12's BF16-accumulating WMMA (the tile takes its step as a template parameter, the served fp32 one by default); gate 4 catches it.
• Tests. Kernel tests against float64 (bound, BF16 rounded once, row bytes at every m and position and on both geometries, exact-sum, cancellation and exponent-boundary probes). hip-dense-tune times both tiles at the served shapes. The qwen3 GPU checks take TF_QWEN3_MATRIX, and projection_test.zig compares both tiles on layer 0's real operands, which layerInputs copies out of the forward (Taps, null in serving).
• Docs. The embedding recipe (the matrix=on serve line beside its tables, gfx1200 unmeasured), docs/kernel-contract.md, and the CHANGELOG.

Checks

• zig build test -Dhip=true: 518/518 steps. Lean: no new findings (424, the base's).
• hip-kernel-test 125/127 (2 skipped); hip-kernel-contract on gfx1201: 64 kernels pass gates 1-4, 6 controls red as planted (the WMMA one among them); the 2 limit problems are upstream's tf_sum_partials records, as on the base. hip-kernel-contract-limits: the WMMA and control records current; its 11 problems are tf_sum_partials, tf_sum_partials_act and tf_moe_combine_gate, stale on the base since upstream's 74f1625.
• qwen3-hip-gpu-test under matrix=on and off: the embedding and reranker checks and the projection comparison pass (the WMMA numbers above equal under on).

[ashhart]

Thanks @laudney. This is a clear, careful PR, and the per-projection check on real operands is exactly what a sum-order
change needs.

Everything checks out on our side:

• It applies on current main once Qwen3-Embedding-8B (BF16) on native HIP: BF16 dense GEMM and a qwen3 embedding engine #1978 is in.
• The tile is gated to gfx12 and to qwen3's BF16 products.
• The bounds and edges are right.
• Commit 23 sets the embedding gate the way we asked.

On the decision you raised: please make WMMA opt-in. Have Kernels.serving pick it only for matrix=on, and keep dot2
under auto. Your numbers are good and both oracles pass. But the fp32 sums are 8 to 41 times further from float64,
and our rule is that a change that makes default output less exact needs an opt-in. Please put the recipe's
--policy matrix=on line next to your speed and accuracy tables, so users can take the 2.5× knowingly.

Two smaller asks:

• a planted control for the WMMA tile in the HIP contract;
• gfx1200's launch record, or a note in the recipe that gfx1200 is unmeasured.

[laudney]

Thanks @ashhart. All three asks are in, as two new commits on top (no force-push):

Ask	Commit	What changed
WMMA opt-in	c5b36f4	Kernels.serving picks WMMA only for matrix=on; auto keeps the dot2 tile and its bits on gfx12. The recipe's --policy matrix=on serve line sits right above its speed and accuracy table, and the CHANGELOG states the opt-in.
Planted control	775339d	tf_ctl_wmma12_bf16_sum: the small block with each 16-deep step through gfx12's BF16-accumulating WMMA (v_wmma_bf16_16x16x16_bf16). The tile takes its step as a template parameter, the served fp32 one by default, so the served kernels' records stay current (80 / 144 VGPRs). On gfx1201: CONTROL control: WMMA dense tile summed in bf16: red at sums, as planted (8.8e10 times over its exact-sum bound at k 12,288).
gfx1200	c5b36f4	No gfx1200 card here, and the limits table has no gfx1200 column, so the recipe and the CHANGELOG say gfx1200 builds the tile but is unmeasured.

Checks after the change: hip-kernel-contract on gfx1201, 64 kernels pass gates 1-4 and 6 of 6 controls red (the 2 limit problems are the base's tf_sum_partials records); qwen3-hip-gpu-test under matrix=on gives the same embedding distances as the WMMA rows in the recipe, the reranker cases pass, and the projection comparison keeps ratio 1.0000 on BF16 out; zig build test -Dhip=true 518/518; lean 424, the base's. The PR description is updated to match.

[ashhart]

Thanks @laudney, all three asks are in.

• Kernels.serving now picks WMMA only for matrix=on, and the host test pins auto to dot2.
• The recipe puts the --policy matrix=on line right above its numbers.
• The BF16-sum control is a clean way to show gate 4 catches a less exact sum.
• The gfx1200 note is clear.

This is ready from our side. It lands after #1978, which needs a small rebase: main has since added the glm_ocr HIP module in zig/build/hip.zig, right where #1978 adds qwen3. When we land it we may fold the two new commits into the ones they amend.

[laudney]

Restacked on #1978's rebased head a5c4e9dee (on main f23603801): the 8 commits replay with no conflict and the same content file by file. On the restacked head: zig build test -Dhip=true 530/530, lean 425 = the new base's, and qwen3-hip-gpu-test under matrix=on passes with the same numbers as before (embedding worst 1.307e-4 at 32,759 tokens; projections BF16 out ratio 1.0000).

[ashhart]

Thank you @laudney. This is merged into main with your commits under your own name, so you're listed as a TensorFold contributor. It ships in the next release, credited to you in the changelog and the release notes.