perf: strix-hip BF16 GEMMs via hipBLASLt (hip-p0→C5 train) by AzeezIsh · Pull Request #166 · Atlas-Inf/atlas

AzeezIsh · 2026-10-06T17:49:39Z

Summary

The cumulative Strix-Halo HIP performance train — hip-p0 (main + #152 + #153) → hip-p2 (+C1 pitched D2D, +C3 w4a16 router twin) → hip-p4 (+C2 FP8-attention skip, +#155, +C4 MoE F16 twins) → C5: route native-HIP BF16 projection GEMMs through hipBLASLt.

Validation on the passing piece (winbox p8/p8b, plan win_plan_1005p8/win_plan_1006p8b):

W1 — Flash-Next prefill A/B vs P4 (3 rotated rounds, greedy, 0 faults): TTFT P5/P4 = 0.899 / 0.903 / 0.927 at 10k/20k/29k. hipBLASLt live in every arm (lt-live=1, lt-warn=0).
W3 — ST-995 (bfcl-subset, n=995, seed 42, temp 0) on P5: overall 84.32 / normalized 84.72 (record run-1791294506763624800) vs P4 refs 84.72/85.38 — within noise, 0 faults, qsa mismatch=0.
Greedy outputs identical at short ctx; expected drift ≥8k (hipBLASLt is not bit-identical to the prior BF16 path).

Strix-27B side (earlier legs): job 161 C5 A/B showed −1.6…2.9% TTFT at 29k on the 27B (marginal — 27B prefill is ~63% w4a16/w8a16 GEMM per the rocprof attribution in job 185), job 168 ST-995 and job 169 replay on the train were clean.

Note: this tip stacks #152, #153, #155 — those land-in-order merges collapse the diff; nothing here changes the default path on non-HIP targets.

🤖 Generated with Devin

…layer host sort past the inert bound) Past the 2051-token inert bound every QSA layer selected its top-512 blocks with a D2H of the scores, a host sort and an H2D upload: 12 stream drains per decode token on Flash-Next. The device arm (`qsa_select_topk`) was opt-in and O(complete^2) (rank every block against every other), measured ~71 us per launch at <= 512 blocks and quadratically worse at agentic lengths; past 4096 blocks the host arm took over regardless. qsa_select_topk_radix selects the identical set in the identical order: - key = order-preserving u32 of rank_key (NaN -> -inf, -0 -> +0), i.e. the host rank_cmp order; - four 8-bit MSD histogram passes find the exact block_topk-th largest key T; every key > T is taken plus the lowest-index keys == T; - two block-wide prefix scans over per-thread contiguous index chunks emit the selection in ascending order, expanded by ratio, then the tail. One 1024-thread CTA, O(complete), up to 16384 blocks (65,536 tokens). The device arm is now the default (ATLAS_QSA_DEVICE_TOPK=0 restores the host sort); QSA_SELECT_MAX_BLOCKS 4096 -> 16384. The old kernel stays for its test, which now runs both device arms against the host reference. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…he dense M128 and routed-MoE kernels gfx1151 has no f32->bf16 instruction: the BF16 prefill kernels round every dequantized weight in software (5 VALU each), which dominates their hot loop (w4a16_gemm_t_m128: 829 instructions per 32-K step for 32 WMMA; the routed MoE kernels: 688 for 16, an older port that also lacked the m128 pipeline). WMMA runs on the same SIMD VALUs, so that work displaces matrix issue. The *_f16a twins keep the signature, grid and launcher of the kernel they replace, and use f16 WMMA: - f16(e2m1) has a zero low byte for every code: one v_perm per four codes; - times the E4M3 scale (bits<<7 * 256, exact incl. subnormals) with one v_pk_mul_f16 per two weights. e2m1*e4m3 is EXACT in f16 (checked over all 254 non-NaN scale bytes x 16 codes), so the dequant no longer rounds; - fp32 per-tensor / per-expert scale2 applied once in the epilogue; - the bf16 A tile converted to f16 while staged in LDS (exact for |a| <= 65504; saturates beyond); - the MoE twins get the m128 register-prefetch / double-buffered pipeline. Hot loop: 378 (dense) and 293 (MoE) instructions per 32-K step. Wiring: ops::strix_f16 picks the twin under cfg!(atlas_hip) and falls back to the BF16 kernel, so GB10 never issues the lookup. Used by dense_ffn, qwen3_attention, qwen3_ssm (m128) and moe/init (gate+up, down). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…-K block scaling) The FP8-weight prefill GEMM decoded each E4M3 weight, multiplied by its 128x128 block scale and rounded to bf16 in software (~11 VALU per weight on gfx1151). The twin decodes E4M3 to f16 exactly (E4M3 is a subset of f16), accumulates each 128-K block unscaled and folds it in once per block (acc += blk * scale, every 4 K-steps), on a uniform condition of the loop counter so it stays a scalar branch. Hot loop 467 -> ~305 instructions per 32-K step (16 WMMA), the flush included every step. Picked by ops::w8a16_n_m128_kernel under cfg!(atlas_hip). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

A BF16 LM head at num_tokens == 2 fell back to two dense_gemv calls, reading the whole weight once per verify token. Qwen3.8-Flash-Next ships its [248320, 2560] LM head BF16, so every MTP verify step read 1.27 GB twice: 2 x 5.87 ms of a 111 ms step on gfx1151 (ATLAS_TRACE_LAUNCH_SYNC, winbox, 2026-10-02). The FP8 and NVFP4 branches already had batch2 arms. Route it to dense_gemv_bf16_batch2 on HIP only. The strix-hip kernel is bit-identical to two dense_gemv_bf16 calls: same 32-lane walk per row and the same shuffle reduction, with the second token as an independent accumulator. The gb10/strix (NVIDIA) builds use 64 lanes per row plus a shared-memory reduction, which reorders the sum, so the handle stays 0 there and the fallback runs unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…A/B attribution Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

strix job 146 (gfx1151, Linux ROCm 7.13), BF16 vs F16 twin TF/s: 27B attn-like 18.29 / 18.12, out-like 17.41 / 18.20, mlp-like 15.56 / 15.06. Neutral, so it does not ship (neutral opt-in switches do not merge). This reverts commit 398d34c. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…40M) strix job 146 (gfx1151), BF16 vs F16 twin TF/s at M=2048: the twin wins on 27B gate/up (+21%), 27B down (+13%) and Flash-Next GDN qkvz (+15%), and loses where BF16 already reaches ~22 TF/s: Flash-Next attn q (-6%), 27B qkv (-2%), Flash-Next out (-9%). N*K >= 40M separates the six shapes exactly. The shared launcher swaps back to the BF16 base below the cut; ATLAS_F16_M128_MIN_NK overrides it. The routed-MoE twins stay ungated (2.1x gate+up, 1.9x down on Flash-Next). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

strix-hip links the cudart->HIP shim (libcudart_hip_shim.cpp maps cudaMemcpy2DAsync to hipMemcpy2DAsync on Linux and Windows), so the per-row copy_d2d_async fallback is unnecessary there. QSA ingest calls copy_d2d_2d_async with 2048 rows per chunk per attention layer; the SYNC launch trace shows ~138 ms of host time per qsa_block_pool launch from those per-row copies.

w4a16_gemm gathers per-element byte loads across 64 rows and pays two barriers per 16-K step; on gfx1151 the Flash-Next router shape (N=512, K=2560) runs it at ~0.3 TF/s at M=2048. The _s64 twin keeps the identical dequant expression, WMMA fragment construction, ascending k-step accumulation order, grid and block; only the staging changes - 64 K per stage (one barrier pair per 4 k-steps), 16-byte A loads into sA[64][72], and per-thread packed B loads (one uint4 + 2 scale bytes) into sB[64][72]. No WMMA is issued for a k-step with kb >= K so the -0.0/+0.0 edge is preserved. Added to both strix-hip w4a16 modules (qwen3.6-27b and qwen3.6-35b-a3b; flash-next symlinks the latter). MoeLayer's w4a16_gemm handle resolves to the twin under cfg!(atlas_hip) via layers::ops::w4a16_gemm_kernel; the lookup is gated so targets without the symbol issue no lookup and leave no failed row in the boot kernel audit.

Native-HIP (strix-hip) lacks the FP8 prefill GEMM kernels - fp8_gemm* and fp8_fp8_gemm* are inline-PTX, not WMMA-ported - so the q/k/v/o FP8 copies built by predequant_for_prefill were consumed by fallback FP8 launches that no strix-hip target ships fast paths for. With the predequant skipped, cache_skip_qkv.rs, paged_qkv.rs and paged_oproj.rs fall through to their existing nvfp4_t arms (w4a16_gemm_t_m128 / w4a16_gemm_n128). Decode never reads q_fp8..o_fp8 (it keys on the native-FP8 source weights via *_weight.as_fp8()). Mirrors the sibling guard in linear_attn_arms.rs. qwen4_exp attention layers build through the same function, so they inherit the guard.

PR #155 swapped MoeLayer::moe_fused_gate_up_t / moe_grouped_gemm_t to their _f16a twins, but no dispatch site ever launches those handles - production prefill (forward_prefill_routed.rs) goes through the _k64 handles. That is why winbox B2 (with #155) vs B1 measured TTFT 1.01x with byte-identical greedy output: the twins never ran. Init now records whether BOTH twins resolved (ops::moe_f16_prefill_engaged); when they did, the two plain transposed-NVFP4 prefill arms launch them via the existing n128 launchers (moe_w4a16_fused_gate_up_n128 / moe_w4a16_grouped_gemm_ptrtable_n128). The twins have the same signatures as their BF16 bases and a 64-row blockIdx.y M tile with no grid striding, so the launchers pass max_m_tiles, not grid_m_strided. Env opt-in arms (m128, fp4, fp8-down) keep priority; e8m0, cutlass and untransposed fallbacks unchanged; non-HIP untouched.

…ad bf16 batch2)

…outer twin)

…155 + C4 (engage routed-MoE F16 twins)

… (C5) Job 160 (gfx1151, ROCm 7.13): hipBLASLt is 2.3-2.6x faster than dense_gemm_bf16_pipelined on every Flash-Next BF16 shape and ~19x the scalar dense_gemm at the 27B MTP fc shape, at bf16-rounding-level error. Four sites now try hipBLASLt first and fall back to their existing kernel on Err, warn-once per site (hip_bf16_proj_or): the mHC collapse tries it at every shape (the machine-fill gate stays NVIDIA-only), the SSM out_proj and force_bf16 QKVZ arms fall back to dense_gemm_bf16_pipelined, and the 27B MTP drafter prefill fc falls back to scalar dense_gemm. The fallback stays on every site — 9fc8e57 moved this path off cuBLASLt on an older ROCm. cfg(atlas_hip) only; NVIDIA paths unchanged.

… cap) moe/mod.rs sat at exactly 500 on main and the train's moe_f16_prefill field pushed it to 502; the init.rs additions pushed it to 501. Move the three pointer-table types (ExpertPtrTable, Fp8ExpertPtrTable, ExpertPtrSet) into a new types.rs re-exported through the moe module — same pattern as ptr_table_build — and tighten the f16-twin comment by one line. mod.rs 458, init.rs 500. cargo check -p spark-model green on strix-hip.

CI 'typos' job on PR #166 fails on two comment occurrences of 'subtile' (suggests subfile/subtitle/subtle). Hyphenate.

AzeezIsh · 2026-10-09T06:52:41Z

Gate-record progress on head 9316140fd (reiner gb10)

7/10 legs now have records stamped at this head (self-served runs on GB10):

ttft-warm-gate, ttft-cold-gate, ssm-state-poisoning-gate, agentic-webserver,
decode-floor, concurrency-sweep, vision-fidelity ✅
Remaining (queued): bfcl-subset (~5h leg), bfcl-subset-echolp, video-fidelity
— the video leg also needs ffmpeg (installed a static arm64 build) and may
hit a KV-budget shortfall: this build commits ~0.5 GB more GPU than the
recipe's util=0.70 allows (84.1 vs 83.7 GB). If it reproduces that's a real
footprint delta vs the baseline records' build.

strix-hip smoke at this head: PASS (build + serve + inference on
qwen3.8-27b/nvfp4, clean). Windows Flash-Next evidence unchanged from the
earlier W1/W3 C5 legs (~10% prefill win).

(Devin)

9 of 10 gate legs measured on GB10 (reiner) at the PR head: ttft-warm-gate, ttft-cold-gate, decode-floor, concurrency-sweep, ssm-state-poisoning-gate, agentic-webserver, vision-fidelity, bfcl-subset (88.14% overall on n=995), bfcl-subset-echolp (87.65%). video-fidelity has no record: its recipe's util=0.70 leaves ~0.4 GB less KV headroom than this build needs (84.1 GB committed vs 83.7 GB budget) at serve time — a footprint delta worth a follow-up look. Generated with [Devin](https://devin.ai) Co-Authored-By: Devin <158243242+devin-ai-integration[bot]@users.noreply.github.com>

Main moved under the gate records (factory/build.rs, m2_setup.rs, transformer_layer.rs since 9316140) — records must cover head+main per the serialized-landing protocol; suite re-run follows on reiner.

9/10 legs re-measured on reiner (GB10) after merging main into the branch — the earlier records stamped 9316140 were invalidated by main-side factory/ changes (serialized-landing protocol). bfcl-subset 88.14%, echolp 87.65% at this tree. video-fidelity still lacks a record: serve fails its util=0.70 KV budget by ~0.6 GB on this ckpt (84.3 committed vs 83.7 budgeted) — either a real footprint delta on this build or a recipe margin issue; follow-up queued. Generated with [Devin](https://devin.ai) Co-Authored-By: Devin <158243242+devin-ai-integration[bot]@users.noreply.github.com>

AzeezIsh · 2026-10-10T08:40:26Z

video-fidelity gate subject — attribution check (gb10)

The recipe self-serve fails at unsloth/Qwen3.6-27B-NVFP4 with
--gpu-memory-utilization 0.70 — on both heads:

build	commit	result
PR head (`0ad989730`, C5+main)	84.3–84.6 GB committed	fails 83.7 GB budget
main (`1d332498`, wt-main binary)	85.0 GB committed	fails identically

So the ~0.9–1.3 GB shortfall is not a C5-train regression — main's build
commits slightly more. The gate subject's pinned budget has gone marginal
against current builds/this ckpt; a record can't be produced by any current
tree without touching the recipe (util bump or a smaller subject). 9/10
records are committed on the PR head at 0ad989730 — video-fidelity is the
only gap and needs a BENCH.toml/subject decision upstream, not another suite
re-run.

(Devin — jobs gb10-video-gate, gb10-video-main-check)

AzeezIsh · 2026-10-14T16:13:05Z

AMD Linux gate at head 61a2c9153 — strix (AzeezStrix, gfx1151, 64 GB, Linux ROCm 7.13).

Same recipes, harness and box as the P5 legs (strix jobs 168/169 on 35b4db9bc), so the variable is the commit: main 7ffd183f2 merged in (incl. the #148/#122 teardown work), the moe types split and the typo fix. Build: strix-hip qwen3.8-27b nvfp4, --no-default-features --features cuda, binary sha256 7a9f8871b885… (strix job 231, both legs).

leg	P5 `35b4db9bc`	head `61a2c9153`
ST-995 (bfcl-subset golden draw, n=995, seed 42, temp 0; 073 recipe: DFlash2 γ=8, util 0.82, seq 26624, prefix caching on)	84.92 / 76.67 (`run-1791229257033697102`)	85.03 / 76.86 (`run-1791982353084802214`)
1007-turn agentic replay (074 recipe, `online_agentic_2.5h_strix_local.yaml`)	1007/1007 completed, 0 failed, score 0.6185, 10474 s	1007/1007 completed, 0 failed, score 0.6209, 10365 s
serve health	0 faults, 0 KV-exhausted, 0 prefill errors, 0 OOM kills	0 faults, 0 KV-exhausted, 0 prefill errors, 0 OOM kills; hipBLASLt live (0 "unavailable" warnings); replay prefix hits 977 (P5: 977)

ST-995 categories match P5 row-for-row except live_irrelevance 44.32 vs 43.18, a single item.
The replay's all_turns_observed compliance line reads FAIL on both heads (observed 1006 = expected 1006 vs issued 1007). P5's run shows the same harness counting; no_dropped_turns passes.
One run per leg, so these are observations under the fingerprints above, not a cross-validated A/B.

Net: no AMD Linux regression from the merge. The only failing check left here is the GB10 video-fidelity gate record (attribution above).

AzeezIsh and others added 18 commits October 2, 2026 02:43

lm_head: gate the BF16 batch2 arm on ATLAS_NO_LM_HEAD_BATCH_GEMV for …

c8f405b

…A/B attribution Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

lm_head: trim the batch2 arm comment under the 500-line cap

397c6e0

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

train/hip-p0: origin/main + PR #152 (qsa radix topk) + PR #153 (lm_he…

d673edd

…ad bf16 batch2)

train/hip-p2: hip-p0 + C1 (pitched D2D on HIP) + C3 (w4a16_gemm_s64 r…

fa692f8

…outer twin)

train/hip-p4: hip-p2 + C2 (skip FP8 attention predequant on HIP) + PR #…

f5a04da

…155 + C4 (engage routed-MoE F16 twins)

kernels: spell sub-tile — typos gate flags subtile

9316140

CI 'typos' job on PR #166 fails on two comment occurrences of 'subtile' (suggests subfile/subtitle/subtle). Hyphenate.

AzeezIsh and others added 3 commits October 9, 2026 09:03

Merge main into perf/hip-bf16-hipblaslt

0ad9897

Main moved under the gate records (factory/build.rs, m2_setup.rs, transformer_layer.rs since 9316140) — records must cover head+main per the serialized-landing protocol; suite re-run follows on reiner.

github-actions Bot added the needs-rebase Conflicts with main — its green checks are NOT real (GitHub skips merge-commit workflows) label Oct 15, 2026

AzeezIsh merged commit fafb734 into main Oct 27, 2026
25 of 26 checks passed

Navigation Menu

Repository navigation

perf: strix-hip BF16 GEMMs via hipBLASLt (hip-p0→C5 train) - #166

perf: strix-hip BF16 GEMMs via hipBLASLt (hip-p0→C5 train)#166
AzeezIsh merged 21 commits into
mainfrom
perf/hip-bf16-hipblaslt

AzeezIsh commented Oct 6, 2026

Uh oh!

AzeezIsh commented Oct 9, 2026

Uh oh!

AzeezIsh commented Oct 10, 2026

Uh oh!

AzeezIsh commented Oct 14, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

AzeezIsh commented Oct 6, 2026

Summary

Uh oh!

AzeezIsh commented Oct 9, 2026

Uh oh!

AzeezIsh commented Oct 10, 2026

Uh oh!

AzeezIsh commented Oct 14, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant