Skip to content

feat(ops): admit the Q6 fused gate_up shape and evaluate it unfused - #284

Open
bingchengcc wants to merge 5 commits into
Neroued:masterfrom
bingchengcc:feat/qwen3.8-27b-q6-profile
Open

bingchengcc wants to merge 5 commits into
Neroued:masterfrom
bingchengcc:feat/qwen3.8-27b-q6-profile

Conversation

@bingchengcc

@bingchengcc bingchengcc commented Sep 19, 2026 •

Copy link
Copy Markdown

What

mlp/gate and mlp/up are 43% of the text weights, and the registered qwen3_8_27b recipe gives
them Q4 - 4.25 bits per weight, 11.5% relative weight error. Q6 reaches 2.6% error at 6.25 bits per
weight. No artifact could use that mixture: linear_swiglu has no Q6 variant, and the generic Q6
GEMM registered only the two vocabulary shapes and the Vision stem.

Three commits, which can be reviewed independently:

  • feat(ops) - registers n=34816, k=5120 as a per-shape Q6 file following the existing
    convention, and evaluates a Q6 gate/up pair through the FFN's existing materialized
    decomposition instead of the fused op.
  • feat(convert) - the qwen3_8_27b_q6 recipe, the guide section, and a characterization test.
  • perf(ops) - tunes the new shape's dispatch over the hot interval and adds the retained
    performance report.

Why the unfused route

linear_swiglu has fp8/nvfp4/q4/q8 but no q6. The FFN already contains an equivalent
decomposition for MTP, so a Q6 gate/up reuses it rather than adding a fifth fused variant. That is
not a compromise: forcing the fused Q4 route through the same decomposition changes the corpus
perplexity below from 1.524538 to only 1.524618.

Tuning

The shape was first registered with the ladder copied from the vocabulary shape, which suits a
geometry with seven times the row-blocks. Measured through production dispatch at every valid T
from 1 to 128:

T vocabulary ladder tuned change
56 330.528 us 228.096 us -31.0%
64 330.880 us 234.624 us -29.1%
80 332.544 us 253.792 us -23.7%
96 334.496 us 274.496 us -17.9%

docs/maintainer/examples/q6-linear.md is the retained report: the curve at every point, the
priority points and both bulk anchors with the bench's logical byte and FLOP accounting, and one
limitation the tuning did not address. T=8..48 runs at 659..701 GB/s against 1320 GB/s at T=1, so
that interval is bound by the k128 MMA route's weight read rather than by its column capacity.

Evidence

Build: cmake --preset release, then -DBUILD_TESTING=ON; clean, CUDA 13.4 / sm_120a.

Op qualification: tests/ops/linear/test_q6_a16.cpp covers the new shape against the suite's FP32
oracle at the selector's ladder boundaries (1, 4, 5, 6, 7, 8, 9, 16, 17, 24, 25, 32, 33, 48, 49,
50, 128, 129). ninfer_linear_q6_a16_test passes, and was re-run after every ladder change.

Converter: pytest tests/convert - 33 passed, 5 of them new in test_official_recipes.py.

End to end, ninfer-perplexity over a 47,917-token corpus with --context 4096 --stride 2048 --kv-dtype int8:

recipe MLP gate/up perplexity
qwen3_8_27b_nvfp4 NVFP4 1.530857
qwen3_8_27b Q4 1.524538
qwen3_8_27b_q6 Q6 1.520765

The artifact is 19.62 GiB and serves 131,072 tokens of int8 KV with 2.63 GiB still free on a single
RTX 5090. Chat, arithmetic and tool-call smoke checks pass.

Scope notes

  • The routing predicate is a pure extension: mtp == true and every non-Q6 format execute the same
    statements as before, so existing artifacts are unaffected.
  • Only gate/up moves. mlp/down and the output projections were also measured at Q6 and rejected:
    linear_add's materialized decomposition rounds the projection to BF16 before adding the
    residual, which costs +0.0021 perplexity, more than the weight gain returns. The vocabulary
    endpoints stay Q8, which at 8.5 bits per weight already outranks Q6.
  • The tuning measurement ran with an idle resident NInfer server on the same GPU; the bench flushes
    256 MiB of L2 before every sample.

Not included

An unrelated local change (ModelOpt NVFP4 import) is deliberately not part of this branch.

The MLP gate/up pair is 43% of the text weights and is the one matrix a Q4
groupwise assignment underserves: 4.25 bits per weight leave 11.5% relative
weight error where Q6 reaches 2.6% at 6.25 bits per weight. Nothing in the
engine could execute that mixture, because linear_swiglu has no Q6 variant and
the generic Q6 GEMM only registered the two vocabulary shapes and the Vision
stem.

Register the fused gate_up geometry (n=34816, k=5120) in the per-shape Q6
dispatch, and evaluate a Q6 gate/up pair through the FFN's existing materialized
decomposition instead of the fused op. That decomposition is not a compromise:
forcing the fused Q4 route through it changes a 47,917-token corpus perplexity
from 1.524538 to 1.524618.

The routing predicate is a pure extension. Both mtp=true and every non-Q6
format take the same statements as before, so existing artifacts are unaffected.

tests/ops/linear/test_q6_a16.cpp qualifies the new shape against the suite's
FP32 oracle at the ladder boundaries the selector distinguishes (1, 4, 5, 6, 7,
8, 9, 16, 17, 24, 25, 32, 33, 48, 49, 50, 128, 129).
_dense_groupwise takes the MLP gate/up format as a parameter and the new
qwen3_8_27b_q6 recipe passes Q6, so the artifact differs from qwen3_8_27b in
exactly those two parameters. The Q8 vocabulary endpoints stay Q8, which at 8.5
bits per weight already outranks Q6.

Measured with ninfer-perplexity over the same 47,917-token corpus, context and
stride 4096/2048 and int8 KV: 1.520765 for q6-full against 1.524538 for
qwen3_8_27b and 1.530857 for qwen3_8_27b_nvfp4.

tests/convert/test_official_recipes.py characterizes the assignment so a later
change to either recipe has to move both together.
The shape was registered with the ladder copied from the vocabulary shape, which is
wrong for this geometry. That shape keeps a single 128-column tile across 49..128
because N=248320 gives it 3880 row-blocks; this N gives 544, so a 128-column tile is
less than half used between 49 and 63.

Measured through production dispatch, every valid T from 1 to 128 plus the two bulk
anchors, 5 warmup and 50 cold-cache samples per point:

  T      vocabulary ladder   tuned     change
  56     330.528 us          228.096   -31.0%
  64     330.880 us          234.624   -29.1%
  80     332.544 us          253.792   -23.7%
  96     334.496 us          274.496   -17.9%
  128    325.792 us          325.600    -0.1%

Splitting 49..64 and 65..96 into 56/64-column and 80/96-column capacities covers the
seven-draft point 56 and the fifteen-draft points 64, 80 and 96. T=97..128 keeps the
128-column capacity: candidates there measured under 1% and do not justify a route.
Across all 128 measured points the change spans -31.0% to +0.7%, the latter at T=25 and
within the repeat-to-repeat spread of a single point.

T=512 and T=1024 are unchanged at 997.984 us (182.91 TFLOP/s, 87.31% TC) and
1868.510 us (195.38 TFLOP/s, 93.26% TC).

docs/maintainer/examples/q6-linear.md records the curve, the priority points and both
anchors with the bench's own logical byte and FLOP accounting, and notes the one
limitation the tuning did not address: T=8..48 sits at 659..701 GB/s against 1320 GB/s at
T=1, so that interval is bound by the k128 MMA route's weight read rather than by the
column capacity.

ninfer_linear_q6_a16_test passes after every ladder change, including the final one.
Wallawalla47 pushed a commit to Wallawalla47/ninfer-custom that referenced this pull request Sep 19, 2026
…the qwen3_8_27b_q6 recipe, and tune the Q6 34816x5120 dispatch
root added 2 commits September 20, 2026 19:08
The k128 small-block MMA tiles move only 64 code bytes per row and stage, so
their 16-byte cp.async requests walk a 2560-byte row stride and the read
efficiency stays near 39% of nominal bandwidth across T=8..48.

Tile the band with BM=32/BK=256 instead: 128 code bytes per row and stage, and
twice the row blocks. T=8..32 drops 13.8%..23.8% (mean 17.8%) while all 128
measured points stay within -23.8%..+0.7% of the previous ladder.

Deeper cp.async pipelines (2 -> 3 -> 4) change nothing, and the BM=16 variant
is slower at T=48, so T=33..64 keeps r64/k128.
The T=1..7 SIMT path stages the weight planes with Cache::ca, but those planes
are read once and never reused from L1; the lines they pin only evict the
activation rows that are reused across warps.

Route the four SIMT schedules through Cache::cg, matching the r64 MMA family.
The c4 tile that serves T=2..4 drops 1.9%..7.1%, T=5 drops 4.1%, and the
c5..c7 tiles are unchanged; all 128 measured points stay within -7.15%..+0.69%
of the previous ladder with no regression above 1%.
@bingchengcc
bingchengcc force-pushed the feat/qwen3.8-27b-q6-profile branch from 97575ff to 4352bdf Compare September 20, 2026 11:39
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant