From c53be4ce90ee3781dbcabba1cc1ec99c6c0d040d Mon Sep 17 00:00:00 2001 From: bingchengcc <175505560+bingchengcc@users.noreply.github.com> Date: Sat, 19 Sep 2026 18:15:35 +0800 Subject: [PATCH 1/5] feat(ops): admit the Q6 fused gate_up shape and evaluate it unfused 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). --- src/models/qwen3_5/execution/ffn.cpp | 26 ++++++++++++++++------- src/ops/linear/q6/q6_dispatch.cpp | 1 + src/ops/linear/q6/q6_shapes.h | 1 + src/ops/linear/q6/shapes/n34816_k5120.cpp | 17 +++++++++++++++ src/ops/linear/q6/sources.cmake | 1 + tests/ops/linear/test_q6_a16.cpp | 7 ++++++ 6 files changed, 45 insertions(+), 8 deletions(-) create mode 100644 src/ops/linear/q6/shapes/n34816_k5120.cpp diff --git a/src/models/qwen3_5/execution/ffn.cpp b/src/models/qwen3_5/execution/ffn.cpp index 21b6b5f35d..d9dea640cd 100644 --- a/src/models/qwen3_5/execution/ffn.cpp +++ b/src/models/qwen3_5/execution/ffn.cpp @@ -22,7 +22,7 @@ std::size_t ffn_workspace_bytes(const FfnParameters& parameters, std::int32_t fi const auto& gu = p.gate_up.weight; const auto& down = p.down.weight; WorkspaceLayoutBuilder layout; - if (mtp) { + if (mtp || gu.qtype == QType::Q6_G64_FP16) { (void)layout.alloc(DType::BF16, {gu.n, last}); { auto scope = layout.scope(); @@ -30,9 +30,15 @@ std::size_t ffn_workspace_bytes(const FfnParameters& parameters, std::int32_t fi gu.qtype, gu.n, gu.k, p.gate_up.policy, first, last)); } (void)layout.alloc(DType::BF16, {gu.n / 2, last}); - (void)layout.alloc(DType::BF16, {down.n, last}); - (void)layout.alloc_bytes(ops::linear_workspace_capacity_bytes(down.qtype, down.n, down.k, - p.down.policy, first, last)); + if (mtp) { + (void)layout.alloc(DType::BF16, {down.n, last}); + (void)layout.alloc_bytes(ops::linear_workspace_capacity_bytes( + down.qtype, down.n, down.k, p.down.policy, first, last)); + } else { + auto scope = layout.scope(); + (void)layout.alloc_bytes(ops::linear_add_workspace_capacity_bytes( + down.qtype, down.n, down.k, p.down.policy, first, last)); + } } else { (void)layout.alloc(DType::BF16, {gu.n / 2, last}); { @@ -65,7 +71,7 @@ void ffn(const Tensor& hidden, const FfnParameters& parameters, Tensor& residual const auto& p = std::get(parameters); const auto& gu = p.gate_up.weight; const auto& down = p.down.weight; - if (mtp) { + if (mtp || gu.qtype == QType::Q6_G64_FP16) { Tensor gate_up = workspace.alloc(DType::BF16, {gu.n, columns}); { auto call = workspace.scope(); @@ -74,9 +80,13 @@ void ffn(const Tensor& hidden, const FfnParameters& parameters, Tensor& residual Tensor activation = workspace.alloc(DType::BF16, {gu.n / 2, columns}); ops::silu_mul(gate_up.slice(0, 0, gu.n / 2), gate_up.slice(0, gu.n / 2, gu.n / 2), activation, stream); - Tensor delta = workspace.alloc(DType::BF16, {down.n, columns}); - ops::linear(activation, down, delta, p.down.policy, workspace, stream); - ops::residual_add(delta, residual, stream); + if (mtp) { + Tensor delta = workspace.alloc(DType::BF16, {down.n, columns}); + ops::linear(activation, down, delta, p.down.policy, workspace, stream); + ops::residual_add(delta, residual, stream); + return; + } + ops::linear_add(activation, down, residual, p.down.policy, workspace, stream); return; } Tensor activation = workspace.alloc(DType::BF16, {gu.n / 2, columns}); diff --git a/src/ops/linear/q6/q6_dispatch.cpp b/src/ops/linear/q6/q6_dispatch.cpp index 1659f09ef6..f2dbe92b92 100644 --- a/src/ops/linear/q6/q6_dispatch.cpp +++ b/src/ops/linear/q6/q6_dispatch.cpp @@ -12,6 +12,7 @@ struct ShapeEntry { constexpr std::array kShapes{ ShapeEntry{248320, 5120, select_q6_n248320_k5120}, + ShapeEntry{34816, 5120, select_q6_n34816_k5120}, ShapeEntry{248320, 2048, select_q6_n248320_k2048}, ShapeEntry{1152, 1536, select_q6_n1152_k1536}, }; diff --git a/src/ops/linear/q6/q6_shapes.h b/src/ops/linear/q6/q6_shapes.h index b1556223b8..1b38cb5d81 100644 --- a/src/ops/linear/q6/q6_shapes.h +++ b/src/ops/linear/q6/q6_shapes.h @@ -5,6 +5,7 @@ namespace ninfer::ops::detail { [[nodiscard]] Q6Launch select_q6_n248320_k5120(std::int32_t tokens); +[[nodiscard]] Q6Launch select_q6_n34816_k5120(std::int32_t tokens); [[nodiscard]] Q6Launch select_q6_n248320_k2048(std::int32_t tokens); [[nodiscard]] Q6Launch select_q6_n1152_k1536(std::int32_t tokens); diff --git a/src/ops/linear/q6/shapes/n34816_k5120.cpp b/src/ops/linear/q6/shapes/n34816_k5120.cpp new file mode 100644 index 0000000000..786cb489ea --- /dev/null +++ b/src/ops/linear/q6/shapes/n34816_k5120.cpp @@ -0,0 +1,17 @@ +#include "ops/linear/q6/q6_shapes.h" + +namespace ninfer::ops::detail { + +Q6Launch select_q6_n34816_k5120(std::int32_t tokens) { + if (tokens <= 4) return launch_q6_simt_r8_c4; + if (tokens <= 5) return launch_q6_simt_r8_c5; + if (tokens <= 6) return launch_q6_simt_r8_c6; + if (tokens <= 7) return launch_q6_simt_r8_c7; + if (tokens <= 16) return launch_q6_mma_r64_c16_k128; + if (tokens <= 24) return launch_q6_mma_r64_c24_k128; + if (tokens <= 32) return launch_q6_mma_r64_c32_k128; + if (tokens <= 48) return launch_q6_mma_r64_c48_k128; + return launch_q6_mma_r64_c128; +} + +} // namespace ninfer::ops::detail diff --git a/src/ops/linear/q6/sources.cmake b/src/ops/linear/q6/sources.cmake index d26b89e072..de17b92b8d 100644 --- a/src/ops/linear/q6/sources.cmake +++ b/src/ops/linear/q6/sources.cmake @@ -1,6 +1,7 @@ target_sources(ninfer_ops PRIVATE "${CMAKE_CURRENT_LIST_DIR}/q6_dispatch.cpp" "${CMAKE_CURRENT_LIST_DIR}/shapes/n248320_k5120.cpp" + "${CMAKE_CURRENT_LIST_DIR}/shapes/n34816_k5120.cpp" "${CMAKE_CURRENT_LIST_DIR}/shapes/n248320_k2048.cpp" "${CMAKE_CURRENT_LIST_DIR}/shapes/n1152_k1536.cpp" "${CMAKE_CURRENT_LIST_DIR}/q6_rowsplit_gemm_mma.cu" diff --git a/tests/ops/linear/test_q6_a16.cpp b/tests/ops/linear/test_q6_a16.cpp index ba498fc54c..919e04da01 100644 --- a/tests/ops/linear/test_q6_a16.cpp +++ b/tests/ops/linear/test_q6_a16.cpp @@ -22,6 +22,13 @@ int q6_a16_conformance() { failures += run_shape("Q6_A16", ActivationCompute::A16, make_q6_g64_fp16_weight, {248320, 5120, 191U, Comparison::Sampled, false, kN248320K5120}); + constexpr std::array kN34816K5120{ + a16(1), a16(4), a16(5), a16(6), a16(7), a16(8), a16(9), a16(16), a16(17), + a16(24), a16(25), a16(32), a16(33), a16(48), a16(49), a16(50), a16(128), a16(129), + }; + failures += run_shape("Q6_A16", ActivationCompute::A16, make_q6_g64_fp16_weight, + {34816, 5120, 199U, Comparison::Sampled, false, kN34816K5120}); + constexpr std::array kN248320K2048{ a16(1), a16(3), a16(4), a16(5), a16(16), a16(17), a16(24), a16(25), a16(32), a16(33), a16(40), a16(41), a16(48), a16(49), a16(56), a16(57), a16(64), a16(65), From d0f36533d7e2ce305ac1693d3056265a717e44ae Mon Sep 17 00:00:00 2001 From: bingchengcc <175505560+bingchengcc@users.noreply.github.com> Date: Sat, 19 Sep 2026 18:15:35 +0800 Subject: [PATCH 2/5] feat(convert): add the qwen3_8_27b_q6 recipe _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. --- docs/weight-conversion.md | 27 ++++++++ tests/convert/test_official_recipes.py | 95 ++++++++++++++++++++++++++ tools/convert/official_recipes.py | 13 ++-- 3 files changed, 131 insertions(+), 4 deletions(-) create mode 100644 tests/convert/test_official_recipes.py diff --git a/docs/weight-conversion.md b/docs/weight-conversion.md index ddeff3be00..2cb81e2720 100644 --- a/docs/weight-conversion.md +++ b/docs/weight-conversion.md @@ -54,6 +54,7 @@ The built-in recipes are ordinary Python functions in |---|---|---| | `qwen3_6_27b` | Q4/Q5 projections, Q6 vocabulary weights | None | | `qwen3_8_27b` | Q4/Q5 projections, Q8 vocabulary weights | None | +| `qwen3_8_27b_q6` | Q4/Q5 projections, Q6 MLP gate/up, Q8 vocabulary weights | None | | `qwen3_6_35b_a3b` | Q4 experts, Q5/Q6 expert down, Q8 shared/projection weights | None | | `qwen3_6_27b_nvfp4` | Imported NVFP4, selected BF16 projections, Q8 vocabulary weights | `quantized` | | `qwen3_8_27b_nvfp4` | Imported NVFP4/FP8, FP8 embedding generated from BF16 | `quantized` | @@ -62,6 +63,32 @@ These names select conversion choices. Runtime execution is selected from the ar configuration and actual bindings stored in the artifact. `--name` sets the public model name; it does not select kernels. +`qwen3_8_27b_q6` differs from `qwen3_8_27b` in one place: the MLP gate and up projections carry Q6 +instead of Q4. Those two parameters are 43% of the text weights, and at 4.25 bits per weight Q4 +leaves precision unused there (11.5% relative weight error against 2.6% for Q6 at 6.25 bits per +weight). The vocabulary endpoints deliberately stay Q8, which at 8.5 bits per weight already +outranks Q6. + +```bash +python3 -m tools.convert \ + --model /path/to/Qwen3.8-27B \ + --recipe qwen3_8_27b_q6 \ + --components text,vision,mtp \ + --resource chat_template.jinja=tools/chat_templates/qwen3_8.jinja \ + --proposal \ + --name qwen3.8-27b-q6 \ + --out models/qwen3_8_27b_q6.ninfer +``` + +Measured with `ninfer-perplexity` over a 47,917-token corpus at `--context 4096 --stride 2048` +with `--kv-dtype int8`, this artifact scores 1.520765 against 1.524538 for `qwen3_8_27b` and +1.530857 for `qwen3_8_27b_nvfp4`. + +The FFN evaluates a Q6 gate/up pair through its existing materialized decomposition (the generic +`linear` op followed by `silu_mul`) rather than the fused `linear_swiglu` op, which has no Q6 +variant. That decomposition costs nothing measurable: forcing the fused Q4 route through the same +decomposition moves the corpus perplexity from 1.524538 to 1.524618. + For a Qwen3.8-27B NVFP4/FP8 artifact with DFlash2: ```bash diff --git a/tests/convert/test_official_recipes.py b/tests/convert/test_official_recipes.py new file mode 100644 index 0000000000..7841fb3796 --- /dev/null +++ b/tests/convert/test_official_recipes.py @@ -0,0 +1,95 @@ +from __future__ import annotations + +import torch + +from tools.convert.model import Model, Parameter +from tools.convert.official_recipes import RECIPES, qwen3_8_27b, qwen3_8_27b_q6 +from tools.convert.recipe import Recipe +from tools.convert.sources.logical import array_source + +Q4 = "q4_g64_fp16" +Q5 = "q5_g64_fp16" +Q6 = "q6_g64_fp16" +Q8 = "q8_g32_fp16" + +LAYER = "text/layers/0" + + +def _dense_model() -> Model: + model = Model({"text": {"config": {}}}) + names = ( + "text/token_embedding", + "text/output_head", + f"{LAYER}/gdn/query", + f"{LAYER}/gdn/output", + f"{LAYER}/attention/query", + f"{LAYER}/attention/key", + f"{LAYER}/attention/output", + f"{LAYER}/mlp/gate", + f"{LAYER}/mlp/up", + f"{LAYER}/mlp/down", + ) + for name in names: + inputs = () if name in ("text/token_embedding", "text/output_head") else ("input",) + source = array_source(torch.ones((4, 8), dtype=torch.bfloat16), name) + model.add(Parameter(name, (4, 8), source, inputs=inputs)) + return model + + +def _formats(recipe_function) -> dict[str, set[str]]: + model = _dense_model() + recipe = Recipe(model) + recipe_function(model, recipe, {}) + return { + name: {selection.format for selection in selections} + for name, selections in recipe.selections.items() + } + + +def _single(formats: dict[str, set[str]], name: str) -> str: + values = formats[name] + assert len(values) == 1, f"{name} is split across formats {values}" + return next(iter(values)) + + +def test_q6_recipe_is_registered() -> None: + assert RECIPES["qwen3_8_27b_q6"] is qwen3_8_27b_q6 + + +def test_registered_recipe_gives_the_mlp_pair_q4() -> None: + formats = _formats(qwen3_8_27b) + assert _single(formats, f"{LAYER}/mlp/gate") == Q4 + assert _single(formats, f"{LAYER}/mlp/up") == Q4 + assert _single(formats, f"{LAYER}/mlp/down") == Q5 + assert _single(formats, "text/token_embedding") == Q8 + assert _single(formats, "text/output_head") == Q8 + + +def test_q6_recipe_moves_only_the_mlp_pair() -> None: + registered = _formats(qwen3_8_27b) + q6 = _formats(qwen3_8_27b_q6) + + assert set(registered) == set(q6) + moved = {name for name in q6 if q6[name] != registered[name]} + assert moved == {f"{LAYER}/mlp/gate", f"{LAYER}/mlp/up"} + + for name in moved: + assert registered[name] == {Q4} + assert q6[name] == {Q6} + + +def test_q6_recipe_keeps_the_vocabulary_at_q8() -> None: + """Q8 is 8.5 bits per weight, so the vocabulary endpoints already outrank Q6.""" + formats = _formats(qwen3_8_27b_q6) + assert _single(formats, "text/token_embedding") == Q8 + assert _single(formats, "text/output_head") == Q8 + + +def test_q6_recipe_leaves_the_other_projections_alone() -> None: + formats = _formats(qwen3_8_27b_q6) + assert _single(formats, f"{LAYER}/attention/query") == Q4 + assert _single(formats, f"{LAYER}/attention/key") == Q4 + assert _single(formats, f"{LAYER}/gdn/query") == Q4 + assert _single(formats, f"{LAYER}/attention/output") == Q5 + assert _single(formats, f"{LAYER}/gdn/output") == Q5 + assert _single(formats, f"{LAYER}/mlp/down") == Q5 diff --git a/tools/convert/official_recipes.py b/tools/convert/official_recipes.py index 3df434ffeb..80919b813c 100644 --- a/tools/convert/official_recipes.py +++ b/tools/convert/official_recipes.py @@ -54,7 +54,7 @@ def _optional(model, recipe): recipe.share(prefix + "context_" + role, prefix + role) -def _dense_groupwise(model, recipe, vocabulary): +def _dense_groupwise(model, recipe, vocabulary, gate_up=Q4): if "num_experts" in model.config: raise ValueError("this official recipe requires Qwen3.5 Dense mathematics") _optional(model, recipe) @@ -66,14 +66,14 @@ def _dense_groupwise(model, recipe, vocabulary): if name.endswith(("/gdn/a_projection", "/gdn/b_projection")): recipe.separate(name) continue - if name.endswith( + if name.endswith(("/mlp/gate", "/mlp/up")): + format = gate_up + elif name.endswith( ( "/attention/query", "/attention/key", "/gdn/query", "/gdn/key", - "/mlp/gate", - "/mlp/up", ) ): format = Q4 @@ -90,6 +90,10 @@ def qwen3_8_27b(model, recipe, sources): _dense_groupwise(model, recipe, Q8) +def qwen3_8_27b_q6(model, recipe, sources): + _dense_groupwise(model, recipe, Q8, gate_up=Q6) + + def qwen3_6_35b_a3b(model, recipe, sources): if "num_experts" not in model.config: raise ValueError("this official recipe requires Qwen3.5 MoE mathematics") @@ -178,6 +182,7 @@ def qwen3_8_27b_nvfp4(model, recipe, sources): "qwen3_6_27b": qwen3_6_27b, "qwen3_6_27b_nvfp4": qwen3_6_27b_nvfp4, "qwen3_8_27b": qwen3_8_27b, + "qwen3_8_27b_q6": qwen3_8_27b_q6, "qwen3_8_27b_nvfp4": qwen3_8_27b_nvfp4, "qwen3_6_35b_a3b": qwen3_6_35b_a3b, } From 1f714da64b382facaa957295bb3cdd7fc7b8ab07 Mon Sep 17 00:00:00 2001 From: bingchengcc <175505560+bingchengcc@users.noreply.github.com> Date: Sat, 19 Sep 2026 18:38:33 +0800 Subject: [PATCH 3/5] perf(ops): tune the Q6 34816x5120 dispatch and report the curve 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. --- docs/maintainer/examples/q6-linear.md | 143 ++++++++++++++++++++++ docs/maintainer/examples/q6-linear.svg | 73 +++++++++++ docs/maintainer/linear-tuning.md | 3 + src/ops/linear/q6/shapes/n34816_k5120.cpp | 4 + 4 files changed, 223 insertions(+) create mode 100644 docs/maintainer/examples/q6-linear.md create mode 100644 docs/maintainer/examples/q6-linear.svg diff --git a/docs/maintainer/examples/q6-linear.md b/docs/maintainer/examples/q6-linear.md new file mode 100644 index 0000000000..df3dacceab --- /dev/null +++ b/docs/maintainer/examples/q6-linear.md @@ -0,0 +1,143 @@ +# Q6 Linear 性能报告：N=34816，K=5120 + +2026-09-19 测量。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能， +供评估该 shape 和后续调优参考。 + +## 测量对象与条件 + +| 项目 | 本次测量 | +|---|---| +| 权重 | `Q6_G64_FP16`，row-split layout；每 64 个权重共享一个 FP16 scale，另有 16 字节高位平面 | +| 数学形状 | `W[34816,5120] × X[5120,T] → Y[34816,T]` | +| 输入、输出与 policy | BF16 输入、BF16 输出，`A16Only` | +| GPU / 工具链 | NVIDIA GeForce RTX 5090；CUDA 13.4，Release，`sm_120a` | +| 计时入口 | `ninfer::ops::linear`；每个 CUDA Graph 包含一次完整 Op 调用 | +| cache / 采样 | 每个样本前清除 256 MiB L2；5 次 warmup、50 次测量，报告 median | +| 输入 fixture | 使用公开 bench 的 Q6 packed-weight 与 BF16 activation fixture | +| 覆盖 | T=1～128 每个整数；512、1024 两个大 T 锚点 | +| 并发条件 | 测量时同机有一个空闲的常驻 NInfer 服务占用显存；bench 每样本清 256 MiB L2 | + +当前被测调用均只有一个 Graph kernel 节点，外部 workspace 为零。测量范围是纯 Linear Op。 + +## 最终耗时曲线 + +![Q6 Linear 最终耗时曲线](q6-linear.svg) + +左图包含全部 128 个实测点，圆点标出重点 T，橙色强调 T=1、4、8。 +右图分别展示 512、1024。图中连线仅连接相邻实测点。 + +## 逻辑带宽与 Tensor Core 利用率 + +权重包含 `34816 × 5120 / 64 × 50 = 139,264,000` 字节（32 字节低位码 + 16 字节高位平面 + +2 字节 FP16 scale，每 64 个权重一组）。沿用 +[Linear bench 指标定义](../linear-benchmark.md#5-数学工作量与-route-neutral-指标)： + +```text +model_bytes = 139,264,000 + 2 × 5120 × T + 2 × 34816 × T +逻辑带宽 GB/s = model_bytes / seconds / 1e9 +带宽利用率 = 逻辑带宽 / 1792 GB/s × 100% + +有效算力 TFLOP/s = 2 × 34816 × 5120 × T / seconds / 1e12 +TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100% +``` + +分母取自 bench 的 RTX 5090 固定规格常量。T=1 使用 GEMV，TC 利用率填 `—`；T≥2 的最终路径 +均使用 BF16 输入、FP32 累加的 dense MMA，使用对应的 **209.5 TFLOP/s** 峰值。 +当前 bench 的 Q6 行未填充 TC 字段，本文按上述公式补算；不加入 padding 或额外 tile 工作量。 + +“READ %”列是相对 bench 另列的 1674.5 GB/s 实测持续读上限的口径。 + +| T | 延迟 µs | 逻辑带宽 GB/s | 带宽利用率 | READ % | 有效算力 TFLOP/s | TC 利用率 | +|---:|---:|---:|---:|---:|---:|---:| +| 1 | 105.504 | 1320.7 | 73.70% | 78.87% | 3.38 | — | +| 2 | 109.280 | 1275.8 | 71.20% | 76.19% | 6.52 | 3.11% | +| 3 | 115.680 | 1205.9 | 67.30% | 72.02% | 9.25 | 4.41% | +| 4 | 121.472 | 1149.1 | 64.12% | 68.62% | 11.74 | 5.60% | +| 5 | 131.744 | 1060.1 | 59.16% | 63.31% | 13.53 | 6.46% | +| 6 | 142.048 | 983.8 | 54.90% | 58.75% | 15.06 | 7.19% | +| 7 | 152.224 | 918.5 | 51.26% | 54.85% | 16.39 | 7.83% | +| 8 | 199.712 | 700.5 | 39.09% | 41.83% | 14.28 | 6.82% | +| 12 | 200.096 | 700.8 | 39.11% | 41.85% | 21.38 | 10.21% | +| 16 | 200.032 | 702.6 | 39.21% | 41.96% | 28.52 | 13.61% | +| 20 | 203.808 | 691.1 | 38.57% | 41.27% | 34.99 | 16.70% | +| 24 | 203.424 | 694.0 | 38.73% | 41.45% | 42.06 | 20.08% | +| 28 | 207.584 | 681.7 | 38.04% | 40.71% | 48.09 | 22.95% | +| 32 | 206.048 | 688.3 | 38.41% | 41.10% | 55.37 | 26.43% | +| 40 | 219.168 | 650.0 | 36.27% | 38.82% | 65.07 | 31.06% | +| 48 | 217.216 | 658.8 | 36.76% | 39.34% | 78.78 | 37.60% | +| 56 | 228.096 | 630.2 | 35.17% | 37.63% | 87.53 | 41.78% | +| 64 | 234.624 | 615.4 | 34.34% | 36.75% | 97.25 | 46.42% | +| 80 | 253.792 | 573.9 | 32.03% | 34.27% | 112.38 | 53.64% | +| 96 | 274.496 | 535.3 | 29.87% | 31.97% | 124.69 | 59.52% | +| 112 | 334.976 | 442.4 | 24.69% | 26.42% | 119.20 | 56.90% | +| 128 | 325.600 | 459.1 | 25.62% | 27.42% | 140.15 | 66.90% | +| 512 | 997.984 | 180.5 | 10.07% | 10.78% | 182.91 | 87.31% | +| 1024 | 1868.510 | 118.3 | 6.60% | 7.07% | 195.38 | 93.26% | + +## 曲线解读与验证 + +[shape 实现](../../../src/ops/linear/q6/shapes/n34816_k5120.cpp) 按 T 选择 SIMT GEMV、 +`k128` 小块 MMA 和固定列宽 MMA 容量： + +- **T=1 为 105.504µs、1320.7 GB/s、 + 标称带宽利用率 73.70%**（持续读口径 + 78.87%）。权重 139.3 MB 构成该点的全部流量。 +- **T=1～7 使用 SIMT 路径**（容量 4、5、6、7），**T=8 起进入 MMA**。T=7→8 是本曲线最大的 + 相邻台阶，`delta_pct` 为 +31.2%，从 152.224µs 到 + 199.712µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 200µs 的固定 + 权重量，两者组织不同，当前保留更低延迟的单 token 专用路径。 +- **T=8～48 使用 `k128` 小块 MMA**，容量 16、24、32、48。跨容量时曲线平滑， + T=16～48 的相邻涨幅均小于 2%。这一段是**带宽受限**的：逻辑带宽 701..659 GB/s、 + 标称利用率 39.1..36.8%，明显低于 T=1 的 73.7%。 + 该区间的耗时几乎与 T 无关（T=8 为 199.712µs、T=48 为 + 217.216µs），说明固定成本是权重读取，尚未达到 T=1 路径的读效率； + 这是当前实现的已知限制，未在本次分派调优范围内。 +- **T=49～128 使用固定列宽容量**，是本形状相对注册词表形状最主要的差别。词表形状 + （N=248320）在 49～128 整段只保留一个 128 列 tile，而本形状的 N 小 7 倍， + `r64` tile 的行块数为 544 而非 3880。按调优规程的候选流程划分后： + +| T | 词表形状的阶梯 | 本形状最终阶梯 | 变化 | +|---:|---:|---:|---:| +| 48 | 217.408 | 217.216 | -0.1% | +| 56 | 330.528 | 228.096 | -31.0% | +| 64 | 330.880 | 234.624 | -29.1% | +| 80 | 332.544 | 253.792 | -23.7% | +| 96 | 334.496 | 274.496 | -17.9% | +| 112 | 335.104 | 334.976 | -0.0% | +| 128 | 325.792 | 325.600 | -0.1% | + + 全部 128 个实测点相对该基线的变化区间为 **-31.0% .. +0.7%**：最大降幅 -31.0%（T=56）， + 最大涨幅 +0.7%（T=25，落在逐点重复测量的波动范围内）。 +- **T=112 到 128 不再细分**：T=112 为 334.976µs，T=128 为 + 325.600µs。候选在 97～128 段测得的收益不足 1%， + 不值得再加一条路由。 +- **大 T 锚点**：T=512 为 **997.984µs、 + 182.91 TFLOP/s、TC 利用率 + 87.31%**；T=1024 为 + **1868.510µs、195.38 TFLOP/s、 + TC 利用率 93.26%**。 + 两者单独测量，避免平均值掩盖其中一个的回归。 + +`ninfer_linear_q6_a16_test` 已通过。新增的 `N=34816, K=5120` case 覆盖 ladder 区分出的 +全部边界（1、4、5、6、7、8、9、16、17、24、25、32、33、48、49、50、128、129），并以现有 +A16 误差标准对独立解码的 packed code 与 FP16 scale 作 FP64 oracle 比较。 +调优过程中每次改动 ladder 都重跑该测试，均通过。 + +## 复现 + +```bash +cmake --build build -j --target ninfer_linear_bench ninfer_linear_q6_a16_test +./build/tests/ninfer_linear_q6_a16_test + +./build/bench/ninfer_linear_bench \ + --qtype q6 --policy a16 --n 34816 --k 5120 \ + --sweep 1:128 --execution graph --warmup 5 --repeat 50 --flush-mib 256 \ + --csv-out profiles/bench/q6_n34816_k5120/final_t1_128.csv +./build/bench/ninfer_linear_bench \ + --qtype q6 --policy a16 --n 34816 --k 5120 \ + --sweep 512:1024:512 --execution graph --warmup 5 --repeat 50 --flush-mib 256 \ + --csv-out profiles/bench/q6_n34816_k5120/final_bulk.csv +``` + +继续调优时，按 [Linear 调优与报告规范](../linear-tuning.md) 选择候选、验证数值并收敛分派； +本报告提供最终实现的测量结果与曲线，不承诺这些容量适用于其他 shape。 diff --git a/docs/maintainer/examples/q6-linear.svg b/docs/maintainer/examples/q6-linear.svg new file mode 100644 index 0000000000..88aa225b72 --- /dev/null +++ b/docs/maintainer/examples/q6-linear.svg @@ -0,0 +1,73 @@ + + + +0 + +91 + +182 + +273 + +364 + +1 + +16 + +32 + +48 + +64 + +80 + +96 + +112 + +128 + + +T (token columns), 1–128 +median latency (µs) +Q6 G64_FP16, N=34816, K=5120 — T=1..128 + + + + + + + + + + + + + + + + + + + + + + + +T=1 +T=4 +T=8 +Bulk anchors + +998.0 µs +183 TFLOP/s +T=512 + +1868.5 µs +195 TFLOP/s +T=1024 + +512 and 1024, measured separately + diff --git a/docs/maintainer/linear-tuning.md b/docs/maintainer/linear-tuning.md index 96a8ffbfdc..e252211f90 100644 --- a/docs/maintainer/linear-tuning.md +++ b/docs/maintainer/linear-tuning.md @@ -83,6 +83,9 @@ requirements for the large-extent region, not a mandatory position in the develo A retained Linear performance report describes the final implementation's absolute performance. Use the [Q4 6144×5120 report](examples/q4-linear.md) as a worked example, with this structure: +Reports for shapes registered later live beside it; the +[Q6 34816×5120 report](examples/q6-linear.md) records the fused gate/up projection at Q6. + 1. State the format/layout, N/K, input/output types, activation policy, GPU/toolchain, timing boundary, cache conditions, warmup/repetitions, and latency statistic. 2. Plot the final latency at every valid T through 128 on linear axes, marking priority points diff --git a/src/ops/linear/q6/shapes/n34816_k5120.cpp b/src/ops/linear/q6/shapes/n34816_k5120.cpp index 786cb489ea..e233a08a8d 100644 --- a/src/ops/linear/q6/shapes/n34816_k5120.cpp +++ b/src/ops/linear/q6/shapes/n34816_k5120.cpp @@ -11,6 +11,10 @@ Q6Launch select_q6_n34816_k5120(std::int32_t tokens) { if (tokens <= 24) return launch_q6_mma_r64_c24_k128; if (tokens <= 32) return launch_q6_mma_r64_c32_k128; if (tokens <= 48) return launch_q6_mma_r64_c48_k128; + if (tokens <= 56) return launch_q6_mma_r64_c56_k128; + if (tokens <= 64) return launch_q6_mma_r64_c64_k128; + if (tokens <= 80) return launch_q6_mma_r64_c80; + if (tokens <= 96) return launch_q6_mma_r64_c96; return launch_q6_mma_r64_c128; } From 19ec27093809b5a9b52f1bebd1090d0cef54aa9c Mon Sep 17 00:00:00 2001 From: root Date: Sun, 20 Sep 2026 19:08:46 +0800 Subject: [PATCH 4/5] perf(ops): route the Q6 8..32 column band to BM=32 BK=256 tiles 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. --- docs/maintainer/examples/q6-linear.md | 129 ++++++++++++++-------- docs/maintainer/examples/q6-linear.svg | 52 ++++----- src/ops/linear/q6/q6_launch.h | 6 + src/ops/linear/q6/q6_rowsplit_gemm_mma.cu | 24 ++++ src/ops/linear/q6/shapes/n34816_k5120.cpp | 6 +- 5 files changed, 144 insertions(+), 73 deletions(-) diff --git a/docs/maintainer/examples/q6-linear.md b/docs/maintainer/examples/q6-linear.md index df3dacceab..26b83b3c62 100644 --- a/docs/maintainer/examples/q6-linear.md +++ b/docs/maintainer/examples/q6-linear.md @@ -1,6 +1,6 @@ # Q6 Linear 性能报告：N=34816，K=5120 -2026-09-19 测量。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能， +2026-09-19 首次测量，2026-09-20 更新 T=8～32 的分派（见文末第二轮调优）。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能， 供评估该 shape 和后续调优参考。 ## 测量对象与条件 @@ -49,28 +49,28 @@ TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100% | T | 延迟 µs | 逻辑带宽 GB/s | 带宽利用率 | READ % | 有效算力 TFLOP/s | TC 利用率 | |---:|---:|---:|---:|---:|---:|---:| -| 1 | 105.504 | 1320.7 | 73.70% | 78.87% | 3.38 | — | -| 2 | 109.280 | 1275.8 | 71.20% | 76.19% | 6.52 | 3.11% | -| 3 | 115.680 | 1205.9 | 67.30% | 72.02% | 9.25 | 4.41% | -| 4 | 121.472 | 1149.1 | 64.12% | 68.62% | 11.74 | 5.60% | -| 5 | 131.744 | 1060.1 | 59.16% | 63.31% | 13.53 | 6.46% | -| 6 | 142.048 | 983.8 | 54.90% | 58.75% | 15.06 | 7.19% | -| 7 | 152.224 | 918.5 | 51.26% | 54.85% | 16.39 | 7.83% | -| 8 | 199.712 | 700.5 | 39.09% | 41.83% | 14.28 | 6.82% | -| 12 | 200.096 | 700.8 | 39.11% | 41.85% | 21.38 | 10.21% | -| 16 | 200.032 | 702.6 | 39.21% | 41.96% | 28.52 | 13.61% | -| 20 | 203.808 | 691.1 | 38.57% | 41.27% | 34.99 | 16.70% | -| 24 | 203.424 | 694.0 | 38.73% | 41.45% | 42.06 | 20.08% | -| 28 | 207.584 | 681.7 | 38.04% | 40.71% | 48.09 | 22.95% | -| 32 | 206.048 | 688.3 | 38.41% | 41.10% | 55.37 | 26.43% | -| 40 | 219.168 | 650.0 | 36.27% | 38.82% | 65.07 | 31.06% | -| 48 | 217.216 | 658.8 | 36.76% | 39.34% | 78.78 | 37.60% | -| 56 | 228.096 | 630.2 | 35.17% | 37.63% | 87.53 | 41.78% | -| 64 | 234.624 | 615.4 | 34.34% | 36.75% | 97.25 | 46.42% | -| 80 | 253.792 | 573.9 | 32.03% | 34.27% | 112.38 | 53.64% | -| 96 | 274.496 | 535.3 | 29.87% | 31.97% | 124.69 | 59.52% | -| 112 | 334.976 | 442.4 | 24.69% | 26.42% | 119.20 | 56.90% | -| 128 | 325.600 | 459.1 | 25.62% | 27.42% | 140.15 | 66.90% | +| 1 | 105.184 | 1324.8 | 73.93% | 79.11% | 3.39 | — | +| 2 | 107.648 | 1295.2 | 72.28% | 77.35% | 6.62 | 3.16% | +| 3 | 115.392 | 1209.0 | 67.46% | 72.20% | 9.27 | 4.42% | +| 4 | 119.648 | 1166.6 | 65.10% | 69.67% | 11.92 | 5.69% | +| 5 | 130.016 | 1074.2 | 59.94% | 64.15% | 13.71 | 6.54% | +| 6 | 141.888 | 984.9 | 54.96% | 58.82% | 15.08 | 7.20% | +| 7 | 151.360 | 923.8 | 51.55% | 55.17% | 16.49 | 7.87% | +| 8 | 166.400 | 840.8 | 46.92% | 50.21% | 17.14 | 8.18% | +| 12 | 169.184 | 828.8 | 46.25% | 49.50% | 25.29 | 12.07% | +| 16 | 168.544 | 833.9 | 46.53% | 49.80% | 33.84 | 16.15% | +| 20 | 168.992 | 833.5 | 46.51% | 49.78% | 42.19 | 20.14% | +| 24 | 168.640 | 837.2 | 46.72% | 50.00% | 50.74 | 24.22% | +| 28 | 158.464 | 893.0 | 49.83% | 53.33% | 62.99 | 30.07% | +| 32 | 158.432 | 895.1 | 49.95% | 53.46% | 72.01 | 34.37% | +| 40 | 216.192 | 658.9 | 36.77% | 39.35% | 65.96 | 31.49% | +| 48 | 213.760 | 669.4 | 37.36% | 39.98% | 80.06 | 38.21% | +| 56 | 226.528 | 634.5 | 35.41% | 37.89% | 88.13 | 42.07% | +| 64 | 233.984 | 617.0 | 34.43% | 36.85% | 97.52 | 46.55% | +| 80 | 254.592 | 572.1 | 31.93% | 34.17% | 112.03 | 53.47% | +| 96 | 275.136 | 534.0 | 29.80% | 31.89% | 124.39 | 59.38% | +| 112 | 338.560 | 437.8 | 24.43% | 26.14% | 117.94 | 56.30% | +| 128 | 327.712 | 456.2 | 25.46% | 27.24% | 139.25 | 66.47% | | 512 | 997.984 | 180.5 | 10.07% | 10.78% | 182.91 | 87.31% | | 1024 | 1868.510 | 118.3 | 6.60% | 7.07% | 195.38 | 93.26% | @@ -79,35 +79,37 @@ TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100% [shape 实现](../../../src/ops/linear/q6/shapes/n34816_k5120.cpp) 按 T 选择 SIMT GEMV、 `k128` 小块 MMA 和固定列宽 MMA 容量： -- **T=1 为 105.504µs、1320.7 GB/s、 - 标称带宽利用率 73.70%**（持续读口径 - 78.87%）。权重 139.3 MB 构成该点的全部流量。 +- **T=1 为 105.184µs、1324.8 GB/s、 + 标称带宽利用率 73.93%**（持续读口径 + 79.11%）。权重 139.3 MB 构成该点的全部流量。 - **T=1～7 使用 SIMT 路径**（容量 4、5、6、7），**T=8 起进入 MMA**。T=7→8 是本曲线最大的 - 相邻台阶，`delta_pct` 为 +31.2%，从 152.224µs 到 - 199.712µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 200µs 的固定 + 相邻台阶，`delta_pct` 为 +9.9%，从 151.360µs 到 + 166.400µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 165µs 的固定 权重量，两者组织不同，当前保留更低延迟的单 token 专用路径。 -- **T=8～48 使用 `k128` 小块 MMA**，容量 16、24、32、48。跨容量时曲线平滑， - T=16～48 的相邻涨幅均小于 2%。这一段是**带宽受限**的：逻辑带宽 701..659 GB/s、 - 标称利用率 39.1..36.8%，明显低于 T=1 的 73.7%。 - 该区间的耗时几乎与 T 无关（T=8 为 199.712µs、T=48 为 - 217.216µs），说明固定成本是权重读取，尚未达到 T=1 路径的读效率； - 这是当前实现的已知限制，未在本次分派调优范围内。 +- **T=8～32 使用 `BM=32`、`BK=256` 的 MMA tile**，容量 16、24、32；**T=33～64 仍使用 + `r64` 系列的 `k128` 小块 MMA**。这一段是**带宽受限**的：T=8～32 的逻辑带宽 + 840.8..895.1 GB/s、标称利用率 46.9..50.0%（第二轮调优后）， + T=33～64 为 669.4..617.0 GB/s、37.4..34.4%。 + 后一段的耗时仍然几乎与 T 无关（T=40 为 216.192µs、T=64 为 + 233.984µs），说明固定成本还是权重读取；受静态 48 KiB 共享内存预算限制， + `BK=256` 的 tile 在 T=48 以上放不下（`BM=16` 的变体又因行块并行度不足更慢）， + 这一段是当前实现的已知限制。 - **T=49～128 使用固定列宽容量**，是本形状相对注册词表形状最主要的差别。词表形状 （N=248320）在 49～128 整段只保留一个 128 列 tile，而本形状的 N 小 7 倍， `r64` tile 的行块数为 544 而非 3880。按调优规程的候选流程划分后： | T | 词表形状的阶梯 | 本形状最终阶梯 | 变化 | |---:|---:|---:|---:| -| 48 | 217.408 | 217.216 | -0.1% | -| 56 | 330.528 | 228.096 | -31.0% | -| 64 | 330.880 | 234.624 | -29.1% | -| 80 | 332.544 | 253.792 | -23.7% | -| 96 | 334.496 | 274.496 | -17.9% | -| 112 | 335.104 | 334.976 | -0.0% | -| 128 | 325.792 | 325.600 | -0.1% | - - 全部 128 个实测点相对该基线的变化区间为 **-31.0% .. +0.7%**：最大降幅 -31.0%（T=56）， - 最大涨幅 +0.7%（T=25，落在逐点重复测量的波动范围内）。 +| 48 | 217.408 | 213.760 | -1.7% | +| 56 | 330.528 | 226.528 | -31.5% | +| 64 | 330.880 | 233.984 | -29.3% | +| 80 | 332.544 | 254.592 | -23.4% | +| 96 | 334.496 | 275.136 | -17.7% | +| 112 | 335.104 | 338.560 | +1.0% | +| 128 | 325.792 | 327.712 | +0.6% | + + 全部 128 个实测点相对该基线的变化区间为 **-31.5% .. +1.0%**：最大降幅 -31.5%（T=56）， + 最大涨幅 +1.0%（T=112，落在逐点重复测量的波动范围内）。 - **T=112 到 128 不再细分**：T=112 为 334.976µs，T=128 为 325.600µs。候选在 97～128 段测得的收益不足 1%， 不值得再加一条路由。 @@ -141,3 +143,42 @@ cmake --build build -j --target ninfer_linear_bench ninfer_linear_q6_a16_test 继续调优时，按 [Linear 调优与报告规范](../linear-tuning.md) 选择候选、验证数值并收敛分派； 本报告提供最终实现的测量结果与曲线，不承诺这些容量适用于其他 shape。 + +## 第二轮调优：T=8～32 换用宽 K tile（2026-09-20） + +第一轮之后，T=8～48 整段的耗时几乎与 T 无关（T=8 199.712µs、T=48 217.216µs）， +`k128` 小块 MMA 的固定成本是权重读取。根因是**每个 stage 每行的字节数太少**： +`BK=128` 时一行一个 stage 只有 64 字节码（32 字节低位 + 16 字节高位按行距 2560 字节分布）， +warp 的 16 字节 `cp.async` 请求跨行跳、线利用率低。把 tile 改成 **`BM=32`、`BK=256`** +后一行一个 stage 达到 128 字节码，同时行块数从 544 翻到 1088，读效率明显上升： + +| T | 第一轮阶梯 µs | 第二轮阶梯 µs | 变化 | +|---:|---:|---:|---:| +| 8 | 197.440 | 166.400 | -15.7% | +| 12 | 197.376 | 169.184 | -14.3% | +| 16 | 197.696 | 168.544 | -14.7% | +| 20 | 203.136 | 168.992 | -16.8% | +| 24 | 202.080 | 168.640 | -16.5% | +| 28 | 205.568 | 158.464 | -22.9% | +| 32 | 205.280 | 158.432 | -22.8% | + +T=8～32 段均值 **-17.84%**（区间 +-23.81% .. -13.80%），T=8 的逻辑带宽从 +700.5 GB/s 升到 840.8 GB/s（标称利用率 39.09% → 46.92%）。 +全部 128 点相对第一轮阶梯的均值变化为 -3.59%， +最大降幅 -23.81%，最大涨幅 +0.72%（落在重复测量波动内）， +没有超过 1% 的回归。 + +第一轮的 `T=7→8` 台阶随之从 +31.2% 收到 +9.9%（151.360µs → 166.400µs）。 + +被否掉的候选（同一计时合同）： + +| 假设 | 实测 | 结论 | +|---|---|---| +| `k128` tile 的流水线深度 2 → 3 → 4 | T=8/16/32 三点在 ±1µs 内 | 不是延迟隐藏不足，加深无收益 | +| `BM=16`、`BK=256` 覆盖 T=48～64 | T=48 从 213.8µs 涨到 251.2µs | 行块并行度不足，该段保留 `r64_c48_k128` | + +数值验证：`ninfer_linear_q6_a16_test` 覆盖 ladder 区分出的全部边界 +（1、4、5、6、7、8、9、16、17、24、25、32、33、48、49、50、128、129）， +以 FP64 oracle 比较独立解码的 packed code 与 FP16 scale，通过。 + diff --git a/docs/maintainer/examples/q6-linear.svg b/docs/maintainer/examples/q6-linear.svg index 88aa225b72..faabbccc06 100644 --- a/docs/maintainer/examples/q6-linear.svg +++ b/docs/maintainer/examples/q6-linear.svg @@ -33,32 +33,32 @@ T (token columns), 1–128 median latency (µs) Q6 G64_FP16, N=34816, K=5120 — T=1..128 - - - - - - - - - - - - - - - - - - - - - - - -T=1 -T=4 -T=8 + + + + + + + + + + + + + + + + + + + + + + + +T=1 +T=4 +T=8 Bulk anchors 998.0 µs diff --git a/src/ops/linear/q6/q6_launch.h b/src/ops/linear/q6/q6_launch.h index 8fa303e6d3..e0f905a8c7 100644 --- a/src/ops/linear/q6/q6_launch.h +++ b/src/ops/linear/q6/q6_launch.h @@ -19,6 +19,12 @@ void launch_q6_mma_r64_c32_k128(const Tensor& x, const Weight& w, Tensor& out, c void launch_q6_mma_r64_c40_k128(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); void launch_q6_mma_r64_c48_k128(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); void launch_q6_mma_r64_c56_k128(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); +void launch_q6_mma_r32_c16_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream); +void launch_q6_mma_r32_c24_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream); +void launch_q6_mma_r32_c32_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream); void launch_q6_mma_r64_c64(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); void launch_q6_mma_r64_c64_k128(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); void launch_q6_mma_r64_c72_k128(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream); diff --git a/src/ops/linear/q6/q6_rowsplit_gemm_mma.cu b/src/ops/linear/q6/q6_rowsplit_gemm_mma.cu index 31c42492ed..583ac9e68d 100644 --- a/src/ops/linear/q6/q6_rowsplit_gemm_mma.cu +++ b/src/ops/linear/q6/q6_rowsplit_gemm_mma.cu @@ -28,6 +28,15 @@ using MmaR64C48K128Schedule = using MmaR64C56K128Schedule = Q6RowSplitMmaGemmSchedule<64, 56, 128, 32, 8, 2, 2, Q6FragmentPipeline::Serial, Cache::cg, Cache::cg, Q6ScaleLoad::Pair32, 1>; +using MmaR32C16K256Schedule = + Q6RowSplitMmaGemmSchedule<32, 16, 256, 16, 8, 2, 3, Q6FragmentPipeline::Serial, Cache::cg, + Cache::cg, Q6ScaleLoad::Pair32, 1>; +using MmaR32C24K256Schedule = + Q6RowSplitMmaGemmSchedule<32, 24, 256, 16, 8, 2, 2, Q6FragmentPipeline::Serial, Cache::cg, + Cache::cg, Q6ScaleLoad::Pair32, 1>; +using MmaR32C32K256Schedule = + Q6RowSplitMmaGemmSchedule<32, 32, 256, 16, 8, 2, 2, Q6FragmentPipeline::Serial, Cache::cg, + Cache::cg, Q6ScaleLoad::Pair32, 1>; using MmaR64C64Schedule = Q6RowSplitMmaGemmSchedule<64, 64, 64, 32, 32, 2, 3, Q6FragmentPipeline::PingPong, Cache::ca, Cache::ca, Q6ScaleLoad::Scalar16>; @@ -121,6 +130,21 @@ void launch_q6_mma_r64_c56_k128(const Tensor& x, const Weight& w, Tensor& out, launch_route(x, w, out, stream); } +void launch_q6_mma_r32_c16_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream) { + launch_route(x, w, out, stream); +} + +void launch_q6_mma_r32_c24_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream) { + launch_route(x, w, out, stream); +} + +void launch_q6_mma_r32_c32_k256(const Tensor& x, const Weight& w, Tensor& out, + cudaStream_t stream) { + launch_route(x, w, out, stream); +} + void launch_q6_mma_r64_c64(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream) { launch_route(x, w, out, stream); } diff --git a/src/ops/linear/q6/shapes/n34816_k5120.cpp b/src/ops/linear/q6/shapes/n34816_k5120.cpp index e233a08a8d..b944fd76d9 100644 --- a/src/ops/linear/q6/shapes/n34816_k5120.cpp +++ b/src/ops/linear/q6/shapes/n34816_k5120.cpp @@ -7,9 +7,9 @@ Q6Launch select_q6_n34816_k5120(std::int32_t tokens) { if (tokens <= 5) return launch_q6_simt_r8_c5; if (tokens <= 6) return launch_q6_simt_r8_c6; if (tokens <= 7) return launch_q6_simt_r8_c7; - if (tokens <= 16) return launch_q6_mma_r64_c16_k128; - if (tokens <= 24) return launch_q6_mma_r64_c24_k128; - if (tokens <= 32) return launch_q6_mma_r64_c32_k128; + if (tokens <= 16) return launch_q6_mma_r32_c16_k256; + if (tokens <= 24) return launch_q6_mma_r32_c24_k256; + if (tokens <= 32) return launch_q6_mma_r32_c32_k256; if (tokens <= 48) return launch_q6_mma_r64_c48_k128; if (tokens <= 56) return launch_q6_mma_r64_c56_k128; if (tokens <= 64) return launch_q6_mma_r64_c64_k128; From 4352bdfca2496386403c8fe05cb5e4f3e5720732 Mon Sep 17 00:00:00 2001 From: root Date: Sun, 20 Sep 2026 19:38:21 +0800 Subject: [PATCH 5/5] perf(ops): stream the Q6 SIMT weight planes through L2 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%. --- docs/maintainer/examples/q6-linear.md | 75 ++++++++++++++-------- docs/maintainer/examples/q6-linear.svg | 48 +++++++------- src/ops/linear/q6/q6_rowsplit_gemm_simt.cu | 8 +-- 3 files changed, 78 insertions(+), 53 deletions(-) diff --git a/docs/maintainer/examples/q6-linear.md b/docs/maintainer/examples/q6-linear.md index 26b83b3c62..37d59313f8 100644 --- a/docs/maintainer/examples/q6-linear.md +++ b/docs/maintainer/examples/q6-linear.md @@ -1,6 +1,6 @@ # Q6 Linear 性能报告：N=34816，K=5120 -2026-09-19 首次测量，2026-09-20 更新 T=8～32 的分派（见文末第二轮调优）。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能， +2026-09-19 首次测量，2026-09-20 更新 T=8～32 的分派与 SIMT 权重流策略（见文末两轮调优）。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能， 供评估该 shape 和后续调优参考。 ## 测量对象与条件 @@ -49,28 +49,28 @@ TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100% | T | 延迟 µs | 逻辑带宽 GB/s | 带宽利用率 | READ % | 有效算力 TFLOP/s | TC 利用率 | |---:|---:|---:|---:|---:|---:|---:| -| 1 | 105.184 | 1324.8 | 73.93% | 79.11% | 3.39 | — | -| 2 | 107.648 | 1295.2 | 72.28% | 77.35% | 6.62 | 3.16% | -| 3 | 115.392 | 1209.0 | 67.46% | 72.20% | 9.27 | 4.42% | -| 4 | 119.648 | 1166.6 | 65.10% | 69.67% | 11.92 | 5.69% | -| 5 | 130.016 | 1074.2 | 59.94% | 64.15% | 13.71 | 6.54% | -| 6 | 141.888 | 984.9 | 54.96% | 58.82% | 15.08 | 7.20% | -| 7 | 151.360 | 923.8 | 51.55% | 55.17% | 16.49 | 7.87% | -| 8 | 166.400 | 840.8 | 46.92% | 50.21% | 17.14 | 8.18% | -| 12 | 169.184 | 828.8 | 46.25% | 49.50% | 25.29 | 12.07% | -| 16 | 168.544 | 833.9 | 46.53% | 49.80% | 33.84 | 16.15% | -| 20 | 168.992 | 833.5 | 46.51% | 49.78% | 42.19 | 20.14% | -| 24 | 168.640 | 837.2 | 46.72% | 50.00% | 50.74 | 24.22% | -| 28 | 158.464 | 893.0 | 49.83% | 53.33% | 62.99 | 30.07% | -| 32 | 158.432 | 895.1 | 49.95% | 53.46% | 72.01 | 34.37% | -| 40 | 216.192 | 658.9 | 36.77% | 39.35% | 65.96 | 31.49% | -| 48 | 213.760 | 669.4 | 37.36% | 39.98% | 80.06 | 38.21% | -| 56 | 226.528 | 634.5 | 35.41% | 37.89% | 88.13 | 42.07% | -| 64 | 233.984 | 617.0 | 34.43% | 36.85% | 97.52 | 46.55% | -| 80 | 254.592 | 572.1 | 31.93% | 34.17% | 112.03 | 53.47% | -| 96 | 275.136 | 534.0 | 29.80% | 31.89% | 124.39 | 59.38% | -| 112 | 338.560 | 437.8 | 24.43% | 26.14% | 117.94 | 56.30% | -| 128 | 327.712 | 456.2 | 25.46% | 27.24% | 139.25 | 66.47% | +| 1 | 105.536 | 1320.3 | 73.68% | 78.85% | 3.38 | — | +| 2 | 107.264 | 1299.8 | 72.53% | 77.62% | 6.65 | 3.17% | +| 3 | 107.424 | 1298.6 | 72.47% | 77.55% | 9.96 | 4.75% | +| 4 | 111.392 | 1253.1 | 69.93% | 74.83% | 12.80 | 6.11% | +| 5 | 126.112 | 1107.5 | 61.80% | 66.14% | 14.13 | 6.75% | +| 6 | 140.032 | 997.9 | 55.69% | 59.60% | 15.28 | 7.29% | +| 7 | 150.592 | 928.5 | 51.81% | 55.45% | 16.57 | 7.91% | +| 8 | 166.752 | 839.0 | 46.82% | 50.10% | 17.10 | 8.16% | +| 12 | 170.560 | 822.1 | 45.88% | 49.10% | 25.08 | 11.97% | +| 16 | 168.928 | 832.0 | 46.43% | 49.68% | 33.77 | 16.12% | +| 20 | 169.728 | 829.9 | 46.31% | 49.56% | 42.01 | 20.05% | +| 24 | 168.704 | 836.9 | 46.70% | 49.98% | 50.72 | 24.21% | +| 28 | 158.400 | 893.3 | 49.85% | 53.35% | 63.02 | 30.08% | +| 32 | 158.528 | 894.6 | 49.92% | 53.43% | 71.97 | 34.35% | +| 40 | 218.112 | 653.1 | 36.45% | 39.01% | 65.38 | 31.21% | +| 48 | 215.808 | 663.1 | 37.00% | 39.60% | 79.30 | 37.85% | +| 56 | 227.008 | 633.2 | 35.33% | 37.81% | 87.95 | 41.98% | +| 64 | 232.704 | 620.4 | 34.62% | 37.05% | 98.05 | 46.80% | +| 80 | 252.896 | 575.9 | 32.14% | 34.39% | 112.78 | 53.83% | +| 96 | 273.248 | 537.7 | 30.01% | 32.11% | 125.25 | 59.79% | +| 112 | 336.672 | 440.2 | 24.57% | 26.29% | 118.60 | 56.61% | +| 128 | 328.320 | 455.3 | 25.41% | 27.19% | 138.99 | 66.35% | | 512 | 997.984 | 180.5 | 10.07% | 10.78% | 182.91 | 87.31% | | 1024 | 1868.510 | 118.3 | 6.60% | 7.07% | 195.38 | 93.26% | @@ -83,8 +83,8 @@ TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100% 标称带宽利用率 73.93%**（持续读口径 79.11%）。权重 139.3 MB 构成该点的全部流量。 - **T=1～7 使用 SIMT 路径**（容量 4、5、6、7），**T=8 起进入 MMA**。T=7→8 是本曲线最大的 - 相邻台阶，`delta_pct` 为 +9.9%，从 151.360µs 到 - 166.400µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 165µs 的固定 + 相邻台阶，`delta_pct` 为 +11.2%，从 150.528µs 到 + 167.296µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 165µs 的固定 权重量，两者组织不同，当前保留更低延迟的单 token 专用路径。 - **T=8～32 使用 `BM=32`、`BK=256` 的 MMA tile**，容量 16、24、32；**T=33～64 仍使用 `r64` 系列的 `k128` 小块 MMA**。这一段是**带宽受限**的：T=8～32 的逻辑带宽 @@ -182,3 +182,28 @@ T=8～32 段均值 **-17.84%**（区间 （1、4、5、6、7、8、9、16、17、24、25、32、33、48、49、50、128、129）， 以 FP64 oracle 比较独立解码的 packed code 与 FP16 scale，通过。 +## 第三轮调优：SIMT 权重流走 L2（2026-09-20） + +`T=1～7` 的 SIMT 路径把权重平面用 `cp.async` 直接送进共享内存，走的是 `Cache::ca`（缓存到 L1）。 +权重是一次性流式读，L1 里不会复用，反而挤掉激活的常驻行。改成 `Cache::cg`（只过 L2）后 +`c4` 档明显变快，`c5～c7` 档中性： + +| T | `ca`（µs） | `cg`（µs） | 变化 | +|---:|---:|---:|---:| +| 2 | 109.248 | 107.136 | -1.9% | +| 3 | 115.680 | 107.616 | -7.0% | +| 4 | 120.000 | 111.424 | -7.1% | +| 5 | 131.680 | 126.272 | -4.1% | +| 6 | 142.144 | 140.352 | -1.3% | +| 7 | 152.384 | 150.528 | -1.2% | + +全部 128 点相对上一轮阶梯的均值变化为 -0.17%，最大降幅 -7.15%，最大涨幅 +0.69%， +没有超过 1% 的回归。四处 `Cache::ca` 同时改为 `Cache::cg`，与 `r64` MMA 系列的策略一致。 + +被否掉的候选（同一计时合同，`c4` 档）： + +| 假设 | 实测 | 结论 | +|---|---|---| +| 增大行块（`RowsPerCta` 8 → 16 → 32） | T=4 从 120.0µs 涨到 120.1 / 126.0µs | 行块并行度已够，加大只增固定开销 | +| 加深 `cp.async` 流水线（2 → 3 → 4） | T=4 从 120.0µs 涨到 130.0 / 136.8µs | 不是延迟隐藏不足 | +| 缩小每 stage 组数（16 → 8） | T=4 从 120.0µs 涨到 128.2µs | 每 stage 字节数太少反而更慢 | diff --git a/docs/maintainer/examples/q6-linear.svg b/docs/maintainer/examples/q6-linear.svg index faabbccc06..158629900a 100644 --- a/docs/maintainer/examples/q6-linear.svg +++ b/docs/maintainer/examples/q6-linear.svg @@ -33,32 +33,32 @@ T (token columns), 1–128 median latency (µs) Q6 G64_FP16, N=34816, K=5120 — T=1..128 - - - - - - - - - - - - + + + + + + + + + + + + - - - - - - - - - -T=1 -T=4 -T=8 + + + + + + + + + +T=1 +T=4 +T=8 Bulk anchors 998.0 µs diff --git a/src/ops/linear/q6/q6_rowsplit_gemm_simt.cu b/src/ops/linear/q6/q6_rowsplit_gemm_simt.cu index 0f94648407..742618c8cd 100644 --- a/src/ops/linear/q6/q6_rowsplit_gemm_simt.cu +++ b/src/ops/linear/q6/q6_rowsplit_gemm_simt.cu @@ -10,10 +10,10 @@ namespace ninfer::ops::detail { namespace { -using SimtR8C4Schedule = Q6RowSplitSimtGemmSchedule<8, 4, 16, 2, Cache::ca, 1>; -using SimtR8C5Schedule = Q6RowSplitSimtGemmSchedule<8, 5, 16, 2, Cache::ca, 1>; -using SimtR8C6Schedule = Q6RowSplitSimtGemmSchedule<8, 6, 16, 2, Cache::ca, 1>; -using SimtR8C7Schedule = Q6RowSplitSimtGemmSchedule<8, 7, 16, 2, Cache::ca, 1>; +using SimtR8C4Schedule = Q6RowSplitSimtGemmSchedule<8, 4, 16, 2, Cache::cg, 1>; +using SimtR8C5Schedule = Q6RowSplitSimtGemmSchedule<8, 5, 16, 2, Cache::cg, 1>; +using SimtR8C6Schedule = Q6RowSplitSimtGemmSchedule<8, 6, 16, 2, Cache::cg, 1>; +using SimtR8C7Schedule = Q6RowSplitSimtGemmSchedule<8, 7, 16, 2, Cache::cg, 1>; template void launch_schedule(const Tensor& x, const Weight& w, Tensor& out, cudaStream_t stream) {