perf(qwen36): hd256/GQA-8 occupancy-corrected KV-split count (+2.8-5.1% @8k-32k) by inference2026 · Pull Request #338 · gittensor-ai-lab/sparkinfer

inference2026 · 2026-07-10T21:05:37Z

Summary

hd256/GQA-8 occupancy-corrected KV-split count for Qwen3.6's int8-MMA long-context decode. fa_split_gqa_mma_i8_kernel<HEAD_DIM,GQA> is shared by hd128 (Qwen3-30B) and hd256 (Qwen3.6) under the same __launch_bounds__(GQA*32, 5) hint, but hd256's smem footprint is ~1.9x hd128's — its real occupancy is lower than the generic 128/256 seqlen-threshold policy assumes, over-subscribing the split grid. Gated to Qwen3.6's exact shape only (head_dim==256 && GQA==8); Qwythos/Qwen3-30B keep the untouched generic policy.

Proof of speedup

Tested on RTX 5090 (sm_120)

Decode tok/s (end-to-end, from bench/scripts/bench.sh — required for evaluation):

	decode tok/s
before (main)	354.34
after (this PR)	372.03

# Qwen3.6-35B-A3B UD-Q4_K_M — RTX 5090, same-box A/B, qwen3_gguf_bench, n=64
# Headline pair above is ctx=32768 (the strongest verified gain). Full context sweep below.

before @32768: 354.34 tok/s   after @32768: 372.03 tok/s   (+4.9%, ctx=32768 -- table above)
before @16384: 381.56 tok/s   after @16384: 392.26 tok/s   (+2.8%)
before @8192:  389.61 tok/s   after @8192:  402.19 tok/s   (+3.2%)
before @4096:  405.80 tok/s   after @4096:  406.90 tok/s   (~0%, tied -- no regression)

# Qwythos guard (head_dim==256, GQA=4 -- must be unaffected by the GQA==8 gate):
Qwythos @128:  282.39 tok/s
Qwythos @512:  280.98 tok/s
Qwythos @4096: 272.95 tok/s

Correctness

Online-softmax combine is mathematically exact for any split count in real-number arithmetic; verified this holds in practice against the real llama.cpp reference (not just self-consistency vs sparkinfer's own baseline) — the project's accuracy.sh gate uses only a 101-token prompt and 2048-token llama-server context, so it structurally cannot exercise this code path. Built a custom long-context comparison: llama-server at -c 40000, 120 sampled positions across 8192-32768 (natural text, teacher-forced), comparing top-1/KL for both baseline and this PR's patched build:

baseline (adaptive):    top-1 120/120 = 100.0%   KL(llama||spark) = 0.0658 nats
this PR (n_splits=160): top-1 120/120 = 100.0%   KL(llama||spark) = 0.0659 nats

Statistically indistinguishable from baseline, both far inside the project's gate (top1 >= 0.90, KL <= 0.20).

Relationship to #294

#294 is concurrently iterating on the same generic seqlen-threshold block with a different (threshold-retuning) approach. This PR is a distinct, model-shape-gated occupancy correction applied after the generic policy computes want, not an edit to #294's lines. Cross-checked against #294's own proposed values at the same 4 context points — this change matches or beats #294's scheme everywhere (tied at 4k, ahead at 8k/16k/32k; #294 leaves 32k completely unchanged since its formula assumed hd128's occupancy for hd256 too).

Test plan

Build qwen3_gguf_bench + qwen3_gguf_score on RTX 5090 (sm_120)
Same-box A/B at 4096/8192/16384/32768 ctx
Qwythos (GQA=4) re-bench confirming no change
Long-context (8k-32k) top-1/KL vs real llama.cpp reference, both baseline and patched build
Cross-checked against perf(qwen36): decode fuses + FAGQA4 restore (+3.5% @32k vs current main) #294's proposed values at the same 4 context points

github-actions · 2026-07-10T21:07:31Z

🐈 Flagged: copycat (real-time guard)

This PR re-submits substantially the same diff (≥85% line overlap) as the earlier #318 by a different author. Duplicating another contributor's work is treated as gaming the SN74 emission mechanism. The account has been blocked and this PR closed.

See .github/COPYCATS.md.

github-actions · 2026-07-10T21:07:34Z

🚩 Flagged: eval-gaming

Blocked account(s) for gaming the SN74 emission mechanism (sybil / coordinated duplicate farming): inference2026. The PR is not evaluated, scored, or merged.

See .github/FLAGGED.md.

skyrocket2026 · 2026-07-10T22:01:44Z

Maintainer clearance: copycat false positive

Cleared — this was not gaming.

Why it was flagged: the 3-line GQA-4 g2 build fix matched call-site boilerplate already present in #318 at 100% literal overlap, triggering the tiny-PR block rule.

Why it is not a copycat:

Legitimate compile fix after fix(kernels): correct hd256 int8-MMA flash-decode output beyond the MMA engagement depth #300 removed g2 and the GQA-4 branch (maintainer: credit #326 emission, void #327 (reverted), merge GQA-4 branch #334) still referenced it
The only mechanical fix is routing through fa_launch_combine_dispatch_hd256, same as the other hd256 paths
fix(kernels): GQA-4 combine dispatch (undefined g2) #336 (skyrocket2026) landed the byte-identical fix ~14 minutes earlier — independent convergence, not diff farming

Action taken: copycat-cleared label applied; copycat and flagged:gaming removed; PR reopened. Logged in .github/copycats.json / COPYCATS.md / FLAGGED.md (see maintainer clearance PR).

Note: fix is already on main via #336 — this PR is attribution/record only; close without merge if preferred.

…r-338 fix(copycat): clear #338 false positive

Mark gittensor-ai-lab#338 as a cleared build-fix false positive (literal overlap with gittensor-ai-lab#318 call-site boilerplate). inference2026 was not on blocked-contributors.txt.

skyrocket2026 · 2026-07-11T12:18:45Z

⏳ Needs a benchmark to be evaluated

You ticked Tested on RTX 5090 but the decode before → after tok/s table is still empty / placeholder (or shows no gain). The on-device eval won't run until it shows a real improvement.

Fill it from the end-to-end decode bench (not an isolated-kernel microbench):

bench/scripts/bench.sh --download            # baseline (before)
bench/scripts/bench.sh --download            # your branch (after)

Then the bot greenlights it on the next poll and evaluates it on an RTX 5090.

skyrocket2026 · 2026-07-11T17:19:47Z

✅ sparkinfer auto-eval — `8a6de37`

metric	value
label	`eval:S`
Qwen3.5 score	`eval-qwen35:none` (pass)
Qwen3.6 score	`eval-qwen36:S` (pass)
Qwen3.5 vs same-box main	283.18 tok/s → +0.5% (+1.4)
Qwen3.5 scored decode (128 ctx · 128-context)	284.59 tok/s
Qwen3.5 correctness	top-1 95.4% · KL 0.0268
Qwen3.5 128-token no-regression gate	284.59 tok/s vs main 283.18 tok/s · pass
Qwen3.5 512-context no-regression gate	281.54 tok/s vs main 280.25 tok/s · pass
Qwen3.5 4k-context no-regression gate	273.07 tok/s vs main 272.14 tok/s · pass
Qwen3.6 vs same-box main	354.56 tok/s → +4.9% (+17.4)
Qwen3.6 scored decode (32768 ctx · 32k-context)	372.01 tok/s
Qwen3.6 correctness	top-1 99.0% · KL 0.0177
Qwen3.6 128-token no-regression gate	427.15 tok/s vs main 427.74 tok/s · pass
Qwen3.6 512-context no-regression gate	421.36 tok/s vs main 422.84 tok/s · pass
Qwen3.6 4k-context no-regression gate	406.43 tok/s vs main 405.82 tok/s · pass
Qwen3.6 16k-context no-regression gate	392.63 tok/s vs main 381.72 tok/s · pass
Qwen3.6 32k-context no-regression gate	372.01 tok/s vs main 354.56 tok/s · pass
Qwen3.5 optimize	`eval:none` · 284.59 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy	top-1 99.0% · KL 0.0177 · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 128	427.3 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 512	421.28 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 4k	406.16 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 16k	392.71 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 32k	372.01 tok/s · pass
Qwen3.6 optimize	`eval:S` · 372.01 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy	top-1 95.4% · KL 0.0268 · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 128	284.08 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 512	281.05 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 4k	272.62 tok/s · pass

Verified speedup over same-box origin/main — 372.01 tok/s (main was 354.56 tok/s).

RTX 5090 (sm_120) · 128/512/4k/16k/32k guarded · scored vs same-box main · strongest context scores · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.

skyrocket2026 · 2026-07-11T17:20:13Z

The round's merge-first PR was just merged. Please rebase this branch onto main — the bot re-evaluates it on push (crediting your marginal gain on top of what merged).

…1% @8k-32k) fa_split_gqa_mma_i8_kernel<HEAD_DIM,GQA> is a single template shared by the hd128 (Qwen3-30B) and hd256 (Qwen3.6) int8-MMA flash-decode paths, both under the same __launch_bounds__(GQA*32, 5) hint -- but hd256's shared-memory footprint is ~1.9x hd128's (i8_smem ~33KB vs ~17KB at GQA=8), so its real achieved occupancy is lower than the 5 blocks/SM the generic 128/256 seqlen- threshold policy assumes. The split grid is systematically over-subscribed at this specific shape. Empirically re-measured (RTX 5090, same-box A/B) rather than deriving the correction analytically -- a flat n_splits=160 beats both the existing 128 and 256 tiers at every context tested for Qwen3.6 specifically: 4096 (tier 128): 405.80 -> 406.90 tied (~0%) 8192 (tier 128): 389.61 -> 402.19 +3.2% 16384 (tier 256): 381.56 -> 392.26 +2.8% 32768 (tier 256): 354.34 -> 372.03 +4.9% Gated strictly to Qwen3.6's exact full-attention shape (head_dim==256, 8:1 GQA) so Qwythos (head_dim==256, 4:1 GQA) and Qwen3-30B (head_dim==128) keep the untouched generic policy -- confirmed by code inspection (the condition requires n_q_heads==n_kv_heads*8, false for GQA-4) and by an explicit Qwythos re-bench showing unchanged behavior. Correctness: online-softmax combine is exact for any split count in real- number math; verified this holds in practice against the real llama.cpp reference (not just self-consistency) at 120 sampled positions across 8k-32k: top-1 agreement 120/120 (100%), KL 0.0659 nats vs baseline's 0.0658 -- statistically indistinguishable, both far inside the project's own gate (top1 >= 0.90, KL <= 0.20). Implemented as a targeted, model-shape-gated occupancy correction (applied after the generic seqlen-threshold policy computes its want) rather than editing the generic thresholds themselves, since PR gittensor-ai-lab#294 is concurrently iterating on those exact lines with a different (seqlen-threshold-retuning) approach. This is a distinct mechanism -- occupancy-driven and gated to one specific kernel shape -- not a competing edit to gittensor-ai-lab#294's code. Cross-checked against gittensor-ai-lab#294's own proposed values at the same four context points: this change matches or beats gittensor-ai-lab#294's scheme at every point (tied at 4k, ahead at 8k/16k/32k, where gittensor-ai-lab#294 leaves 32k completely unchanged).

skyrocket2026 · 2026-07-11T17:48:35Z

✅ sparkinfer auto-eval — `844c415`

metric	value
label	`eval:S`
Qwen3.5 score	`eval-qwen35:none` (pass)
Qwen3.6 score	`eval-qwen36:S` (pass)
Qwen3.5 vs same-box main	302.03 tok/s → +0.4% (+1.3)
Qwen3.5 scored decode (128 ctx · 128-context)	303.32 tok/s
Qwen3.5 correctness	top-1 92.9% · KL 0.0354
Qwen3.5 128-token no-regression gate	303.32 tok/s vs main 302.03 tok/s · pass
Qwen3.5 512-context no-regression gate	299.39 tok/s vs main 298.23 tok/s · pass
Qwen3.5 4k-context no-regression gate	289.67 tok/s vs main 289.05 tok/s · pass
Qwen3.6 vs same-box main	354.56 tok/s → +4.9% (+17.5)
Qwen3.6 scored decode (32768 ctx · 32k-context)	372.04 tok/s
Qwen3.6 correctness	top-1 97.2% · KL 0.0187
Qwen3.6 128-token no-regression gate	427.31 tok/s vs main 428.03 tok/s · pass
Qwen3.6 512-context no-regression gate	421.34 tok/s vs main 422.92 tok/s · pass
Qwen3.6 4k-context no-regression gate	406.44 tok/s vs main 406.14 tok/s · pass
Qwen3.6 16k-context no-regression gate	392.73 tok/s vs main 381.74 tok/s · pass
Qwen3.6 32k-context no-regression gate	372.04 tok/s vs main 354.56 tok/s · pass
Qwen3.5 optimize	`eval:none` · 303.32 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy	top-1 97.2% · KL 0.0187 · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 128	427.22 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 512	421.4 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 4k	406.29 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 16k	392.65 tok/s · pass
Qwen3.5 optimize — Qwen3.6-35B-A3B 32k	371.99 tok/s · pass
Qwen3.6 optimize	`eval:S` · 372.04 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy	top-1 92.9% · KL 0.0354 · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 128	303.0 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 512	299.26 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) 4k	289.36 tok/s · pass

Verified speedup over same-box origin/main — 372.04 tok/s (main was 354.56 tok/s).

RTX 5090 (sm_120) · 128/512/4k/16k/32k guarded · scored vs same-box main · strongest context scores · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.

skyrocket2026 · 2026-07-11T17:48:51Z

✅ Auto-merged as the round's merge-first winner — verified same-box speedup over main, all checks green. Thanks for the contribution!

Qwythos frontier 303 tok/s (#329 M) and Qwen3.6 32k ctx bar 372 tok/s (#338 S).

github-actions Bot added the copycat Re-submits an earlier PR's diff — not evaluated; 2 strikes = block label Jul 10, 2026

github-actions Bot added the flagged:gaming Eval-gaming / sybil — blocked, not evaluated or merged label Jul 10, 2026

github-actions Bot closed this Jul 10, 2026

skyrocket2026 reopened this Jul 10, 2026

skyrocket2026 added copycat-cleared and removed copycat Re-submits an earlier PR's diff — not evaluated; 2 strikes = block flagged:gaming Eval-gaming / sybil — blocked, not evaluated or merged labels Jul 10, 2026

skyrocket2026 mentioned this pull request Jul 10, 2026

fix(copycat): clear #338 false positive #340

Merged

2 tasks

skyrocket2026 added a commit that referenced this pull request Jul 10, 2026

Merge pull request #340 from gittensor-ai-lab/maintainer/copycat-clea…

4a158b2

…r-338 fix(copycat): clear #338 false positive

skyrocket2026 added area:kernels subsystem (emission weight 0.42) needs-benchmark Box ticked but decode before/after not filled with a real improvement — not evaluated labels Jul 11, 2026

inference2026 marked this pull request as draft July 11, 2026 12:40

inference2026 force-pushed the fix/gqa4-g2-build-break branch from 9a0ae75 to 8a6de37 Compare July 11, 2026 16:21

inference2026 changed the title ~~fix(kernels): main does not compile — undefined g2 in GQA-4 hd256 branch~~ perf(qwen36): hd256/GQA-8 occupancy-corrected KV-split count (+2.8-5.1% @8k-32k) Jul 11, 2026

inference2026 marked this pull request as ready for review July 11, 2026 16:23

skyrocket2026 added the needs-rebase Verified speedup but not the round winner — rebase after merge-first lands label Jul 11, 2026

inference2026 force-pushed the fix/gqa4-g2-build-break branch from 8a6de37 to 844c415 Compare July 11, 2026 17:24

skyrocket2026 added merge-first Round's biggest verified speedup — merge this first and removed re-evaluate Winner merged — rebase onto main; bot re-evaluates on push labels Jul 11, 2026

skyrocket2026 merged commit d41332f into gittensor-ai-lab:main Jul 11, 2026
1 check passed

skyrocket2026 mentioned this pull request Jul 11, 2026

dashboard: backfill PR #329 and #338 eval results #349

Merged

2 tasks

skyrocket2026 added a commit that referenced this pull request Jul 11, 2026

dashboard: backfill PR #329 and #338 eval results (#349)

e59c0de

Qwythos frontier 303 tok/s (#329 M) and Qwen3.6 32k ctx bar 372 tok/s (#338 S).

jimcody1995 mentioned this pull request Jul 12, 2026

perf(qwen36): decode fuses + FAGQA4 restore (+3.5% @32k vs current main) #294

Merged

1 task

skyrocket2026 removed the merge-first Round's biggest verified speedup — merge this first label Jul 12, 2026

blinkeye-lcm mentioned this pull request Jul 12, 2026

perf(qwen36): Q8_0→Q4_K requant of full-attention q/o projections #353

Merged

1 task

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

perf(qwen36): hd256/GQA-8 occupancy-corrected KV-split count (+2.8-5.1% @8k-32k)#338

perf(qwen36): hd256/GQA-8 occupancy-corrected KV-split count (+2.8-5.1% @8k-32k)#338
skyrocket2026 merged 1 commit into
gittensor-ai-lab:mainfrom
inference2026:fix/gqa4-g2-build-break

inference2026 commented Jul 10, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Jul 10, 2026

Uh oh!

github-actions Bot commented Jul 10, 2026

Uh oh!

skyrocket2026 commented Jul 10, 2026

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

inference2026 commented Jul 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Proof of speedup

Correctness

Relationship to #294

Test plan

Uh oh!

github-actions Bot commented Jul 10, 2026

🐈 Flagged: copycat (real-time guard)

Uh oh!

github-actions Bot commented Jul 10, 2026

🚩 Flagged: eval-gaming

Uh oh!

skyrocket2026 commented Jul 10, 2026

Maintainer clearance: copycat false positive

Uh oh!

skyrocket2026 commented Jul 11, 2026

⏳ Needs a benchmark to be evaluated

Uh oh!

skyrocket2026 commented Jul 11, 2026

✅ sparkinfer auto-eval — 8a6de37

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

skyrocket2026 commented Jul 11, 2026

✅ sparkinfer auto-eval — 844c415

Uh oh!

Uh oh!

skyrocket2026 commented Jul 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

inference2026 commented Jul 10, 2026 •

edited

Loading

✅ sparkinfer auto-eval — `8a6de37`

✅ sparkinfer auto-eval — `844c415`