Skip to content

Speed up native Bolt packed projections and IQ2 experts - #34002

Open
Rui Ren (rui-ren) wants to merge 10 commits into
rui-ren/bolt-standalone-router-on-33969-20261108from
rui-ren/bolt-packed-latency-on-33994-20261109
Open

Rui Ren (rui-ren) wants to merge 10 commits into
rui-ren/bolt-standalone-router-on-33969-20261108from
rui-ren/bolt-packed-latency-on-33994-20261109

Conversation

@rui-ren

@rui-ren Rui Ren (rui-ren) commented Nov 9, 2026 •

Copy link
Copy Markdown
Contributor

Description

Optimizes the existing CUDA MatMulPacked and IQ2_XXS MoEPacked paths without changing
operator schemas, packed-weight layouts, or exported graphs. NVFP4 uses vectorized
group decoding, paired BF16 operands, and half-chunk staging for eligible large-N
Hopper projections. Q8 uses aligned packed-record windows with stride/alignment
fallbacks; dense IQ2 reconstructs complete BF16 operands using the original K partitions.

IQ2 experts use exact magnitude-codebook reconstruction and launch-selected aligned
FC1 activation loads. FP32 scale reconstruction, sign-before-BF16 rounding, FC1/FC2
warp partitions, ReLU-square boundaries, and final reduction order are preserved.
The shared LoadPacked8 fallback from the standalone-router extraction is reused.

Dependency: #33994 in
native stack #33967.
Mamba and standalone/integrated routing are unchanged. No requantization, GenAI
inference change, graph conversion, or acceptance-threshold adjustment.

Summary of Changes

Packed Projection and Expert Kernels

File Change
onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh Add vectorized NVFP4/Q8 loads and dense IQ2 BF16 operand reconstruction.
onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu Select eligible optimized paths while retaining architecture/alignment fallbacks.
onnxruntime/contrib_ops/cuda/moe/moe_packed_iq2.cuh Add IQ2 expert codebook/operand helpers with exact FP32 scale reconstruction.
onnxruntime/contrib_ops/cuda/moe/moe_packed.cu Use the expert helpers and aligned FC1 activation-load specialization.

Regression Tests

File Change
onnxruntime/test/contrib_ops/matmul_packed_precision_test.cc Cover production projection widths, NVFP4 operand extremes, Q8 record windows/stride fallback, and IQ2 complete operands.
onnxruntime/test/contrib_ops/moe_packed_routing_test.cc Add production-width and IQ2 codebook/scale reconstruction cases.
onnxruntime/test/contrib_ops/cuda_kernels/matmul_packed_kernel_test.cc Pin the compiled TC contiguous-128-K reduction for small/staged NVFP4 paths.
onnxruntime/test/python/transformers/test_gguf_packed_ops.py Exercise default-rounded byte-offset IOBinding fallbacks without a model precision attribute.

Testing

Session-policy compatibility update: head 6294c1fd47b493fc5dbbbed6cfad3cc406ca476e, rebased onto
the default-rounded MatMulPacked contract from #33962. The removed node attribute is no
longer emitted by test model factories or assigned through internal launch parameters.
Default numerical references round weights for both activation dtypes and every dispatch.

Validation at optimized stack tip a3f2076b65e0d1a5c7e9b82165ffab014e233876: 15 native CUDA + 15
compute75-only PTX projection tests, 5 Python format/policy/alignment tests, and 1 actual
internal CUDA contiguous-128-K cancellation test; zero failures/skips. CLANGFORMAT,
Ruff/format and whitespace checks pass.

These are focused ABI-matched overlays with fresh changed objects, not a clean full-source
build, full CI, performance remeasurement, or new full-model qualification
. The previous
broader review qualification and historical performance results are preserved below.

Historical measurements and qualification evidence (not current-head acceptance)

Earlier review qualification (historical)

Review-fix snapshot: 74847ad6021d5b5431fa5984efb7fa95f6a310cd.

Added offset-buffer IOBinding fallbacks (NVFP4/Q8 and IQ2 ReLU-square MoE), nonuniform single-product and bitwise 8-vs-9-row MoE oracles, all-256-position IQ2 one-hot coverage, NVFP4 high-nibble coverage and an actual internal contiguous-128-K cancellation regression. Removed duplicate IQ2 table/decoder logic and dead branches; lowered the operand-extremes fixture gate to SM80. Seven wrong-mapping mutations produce nonzero output differences in the new oracle.

Owning-layer validation: 20 C++ + 2 Python IOBinding + 1 actual internal CUDA cancellation test; zero failures/skips.

Integrated tip regression: 51 C++ provider/schema tests, 5 actual internal CUDA tests and 2 Python IOBinding tests passed, zero failures/skips. All owning-layer provider cases are retained. The final overlay recompiles the changed CUDA/host/schema/test sources over ABI-matched preserved ORT support; not a clean full-source build, full-CI approval, independent HF parity or new end-to-end qualification. Targeted configured clang-format/Ruff and whitespace checks passed.

These previously reported results are retained for provenance. They have not been rerun on the rebased head above; earlier CI/build status and benchmark results do not establish current-head CI, accuracy, or merge readiness.

Original comparison pins: parent #33994 at a8ff5642f5c85519561964928006413a4fde7859; candidate 1b5063c4af651e12bf955bc978c9312e2734a9f9.

This is the actual stack-tip comparison, not the earlier private-branch
versus #33940 comparison. That distinction matters: the remote stack tip did
not yet contain the optimized native kernels.

Input tokens Stack ordinary tok/s Candidate ordinary tok/s Stack k=7 tok/s Candidate k=7 tok/s
178-254 54.36 92.42 223.84 355.42
427-932 54.20 92.08 185.24 293.62
1028-1663 54.25 92.04 201.48 318.91

Approximately 70% ordinary and 58-59% k=7 improvement in this controlled
comparison. This does not claim that all remaining latency is resolved.
The separate #33940 equivalent-package result of approximately 109 ordinary /
370 k=7 tok/s on the longest bin is not a measurement of this candidate;
transferring that remaining gain is still outstanding.

Protocol: original frozen 30 real prompts, ten per bin, unchanged warmups and
per-prompt mode order, natural EOS/512-new-token cap, two reversed-order passes,
same physical H100 per baseline/candidate shard. Eight logical shards ran in
waves over GPUs 1,3,4,5,6,7. All admitted runs passed the foreign-client and
resource guards. No profiler ran in the timing window; clocks were not locked.
Rates pool decoded tokens over elapsed decode time, excluding the first token,
not arithmetic means of per-request TPS.

Correctness and build validation

  • Exact same native model package, packed bytes and GenAI libraries in both arms.
    Graph hashes match the preserved source-value audit; full package-file hashes
    were recorded again. No exported checkpoint/package was modified.
  • 22 native CUDA regressions passed, zero skips, covering projections,
    complete operand rounding/cancellation, IQ2 production widths/codebook scales,
    routing and the standalone router. Changed C++ files passed CLANGFORMAT;
    git whitespace checks passed.
  • 1,782 complete live tensor/state hashes per arm, bitwise equal and finite
    in ordinary and k=7 modes, also equal to the preserved quantized control.
  • Independent full LCB145 ordinary/k=7 probes passed all 512 generated tokens
    and complete counters.
  • The accepted timing run passed 240 complete measured generation sequences
    against frozen ordinary-control tokens. Within-mode complete speculative
    counter dictionaries matched between arms and repeats.
  • Isolated incremental rebuild on the archived clean Reuse MoEGroupRouter for projected sigmoid routing #33994 build; every
    affected compilation unit/header dependent was rebuilt and the provider/test
    binary relinked with the production flags. Untouched objects were reused:
    this is not a new clean full-source build.
  • CUDA 12.8.93, driver 580.159.04, H100 80GB. Common corrected GenAI source:
    c0fd01914265d5157d6da1376c38b4df521a76c5
    (microsoft/onnxruntime-genai#2916).
    Both arms use that same source; its correctness fix is not credited with the
    kernel speedup.

Accepted statistics SHA256:
0b9b7e2385b2dd9d7e5f7d459c65ef74338b5f707d1536454e7e27fb9a407c0f.

Checklist

  • Review feedback addressed with focused regressions
  • Documentation and supported contracts updated
  • Native stack order and immediate-parent dependencies preserved
  • Clean full-source/full-CI qualification (not claimed by focused overlays)

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The tests leave NVFP4 pointer-alignment fallbacks unverified; exact-result BF16 cases with complete-operand rounding are needed for unaligned B and the SM90+ unaligned-A staged-load path.

1 open issue
What changed in this PR

This PR optimizes native CUDA packed projections and IQ2 experts on top of #33994, preserving the existing MatMulPacked and MoEPacked interfaces, packed layouts, and routing integration.

Changes:

  • Adds vectorized NVFP4 and Q8 decoding, including staged NVFP4 loads for eligible Hopper projections.
  • Optimizes dense and expert IQ2 operand reconstruction while retaining rounding and reduction boundaries.
  • Expands regression coverage for production shapes, codebook values, numerical extremes, and malformed routes.
File Description
onnxruntime/​test/​contrib_ops/​moe_packed_routing_test.cc Adds IQ2 width/codebook coverage and malformed-route cases.
onnxruntime/​test/​contrib_ops/​matmul_packed_precision_test.cc Adds packed-projection precision and shape regressions.
onnxruntime/​contrib_ops/​cuda/​quantization/​matmul_packed.cu Integrates optimized decoders and alignment-aware dispatch.
onnxruntime/​contrib_ops/​cuda/​quantization/​gguf_quant_cuda.cuh Adds vectorized Q8, NVFP4, and BF16 IQ2 decoding helpers.
onnxruntime/​contrib_ops/​cuda/​moe/​moe_packed.cu Uses optimized IQ2 fragments and aligned FC1 activation loads.
onnxruntime/​contrib_ops/​cuda/​moe/​moe_packed_iq2.cuh Adds expert-specific IQ2 fragment reconstruction.

Review effort: Balanced

@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch from da0f239 to 99f4ab7 Compare November 9, 2026 18:34
rui-ren

This comment was marked as resolved.

@tianleiwu Tianlei Wu (tianleiwu) left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reviewed this layer (stack #33967, layer 5) at 99f4ab7c63.

Kernel review. I extracted the parent (#33994) and PR-head kernels into a standalone harness (H200, sm_90) and ran both on identical random, non-uniform inputs: 2,190 dense MatMulPacked configs (NVFP4/Q8/IQ2, N 1..18,560, K 64..8,192, rows 1..8, B offsets 0/1/2/4/16 B, A offsets 0/2/6/16 B, with and without global_scale, injected 0x7f/0xff scales) and 140 MoE FC1/FC2 configs (aligned and misaligned input, 8- and 32-route tiles) are bit-identical, and so is the pre-SM90 fallback (forced on sm_90). Exhaustive decoder checks (16.8M NVFP4 groups, 1.05M IQ2 sub-blocks) show 0 mismatches against the pre-PR decode. compute-sanitizer memcheck (exact-size allocations), racecheck and synccheck were clean, and clang-format is clean. Kernel timings improve with no regression (e.g. NVFP4 N=18560/K=4096 137 -> 43 us, Q8 103 -> 63 us, IQ2 dense 71 -> 57 us, MoE FC1/FC2 1.1-1.6x). I found no correctness, memory-safety or numerics defect. (The end-to-end numbers in the description were not reproduced.)

Main request: tests that cannot fail for the code this PR rewrites (inline F1-F4). Replaying the tests' own inputs through mutated decoders: 7 of 7 plausible bugs in MoEPackedIq2Fragment, 3 of 3 dense-IQ2 layout bugs and an NVFP4 high-nibble lane swap all pass, and disabling the NVFP4 128-K wide dispatch also passes. #34003 keeps editing the same functions, so this is cheapest to add here.

Existing thread. The unaligned-pointer fallback tests (Copilot thread on matmul_packed_precision_test.cc) are still open and the checklist item in the PR description is unchecked; I added details there (the gap also covers the new Q8 base-pointer clause and the MoE FC1 aligned-input selection).

Other

  • Maintainability: moe_packed_iq2.cuh re-adds a second 256-entry IQ2 BF16 table and helpers that this PR also adds to gguf_quant_cuda.cuh; reusing gguf_cuda::kIq2xxsGridBf16 removes about 290 lines with bit-identical results (inline).
  • Nits (inline): the SM90 gate on Nvfp4Bf16PackedOperandExtremes, and a dead adjacent ternary in the generic decode branch.
  • The NVFP4 tests here pin signed-E4M3 semantics for scale bytes >= 0x80 (if (scale & 128) d = -d, 0xFF decodes to NaN), which differs from the pinned ggml_ue4m3_to_fp32 reference (see my note on #33962). That is fine as a decision, but it should be written into docs/PackedWeightFormats.md.
  • CI: Kernel Documentation Generation fails on this PR; that is inherited from #33994 (the MoERouterPacked docs entry, see my comment there), not caused by this layer.

@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch from 99f4ab7 to 74847ad Compare November 9, 2026 23:43
@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch 2 times, most recently from 79f5c64 to 6852d7f Compare November 10, 2026 00:58

@tianleiwu Tianlei Wu (tianleiwu) left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Re-review of 6852d7fdc4.

Verdict: COMMENT — no defects found in this layer's code; it is red only because of #33994's compile errors. Two non-blocking suggestions and one robustness nit are inline.

Earlier threads, re-checked on this head — all fixed:

  • alignment fallbacks are now reachable (Python IOBinding tests with byte offsets: NVFP4/Q8 A +2 / B +1, and the ReLU² MoE input);
  • nonuniform MoE IQ2 coverage (single-product oracle, bitwise 8-vs-9-row equivalence, random data against the CPU dequantizer);
  • one-hot sweep over all 256 IQ2 block positions; NVFP4 in-group positions 8–14 are activated;
  • the internal MatMulPackedNvfp4WideContiguous128KCancellation pins the contiguous 128-K partition;
  • the duplicated IQ2 table/decoder is gone (shared kIq2xxsGridBf16 / DecodeIq2SubBlock16Bf16); the NVFP4 tests are gated on SM80.

Independent verification (stack tip 623a725f5c, which contains this layer; Linux, CUDA 13.0, sm_80;sm_90a, Release; A100-SXM4-80GB)

  • Build OK; the stack's 60 provider tests and 5 internal CUDA tests pass (also 3x shuffled); the pre-existing *Packed* suites (655 tests) show 0 failures; compute-sanitizer synccheck, racecheck and initcheck are clean on the MatMulPacked / MoEPacked tests.
  • Standalone harness compiled for sm_80 around the unmodified device helpers:
    • DecodeNvfp4WordsBf16 (the Hopper BF16 operand reconstruction) vs. the FP32 formula d * (0.5f * kvalues_fp4[code]) rounded to BF16, for every scale byte 0x00–0xFF × all 16 nibble values at all 16 positions + 4096 random 64-bit code words per scale: 0 mismatches in 591,872 in-contract decodes (and 0 in as many out-of-contract ones);
    • DecodeIq2SubBlock16Bf16 (all four kCodesWord / kSignBits variants, and the aligned-record-window variant on even block counts) and DecodeQ8SubBlock16Aligned vs. the CPU dequantizer rounded to BF16: 0 mismatches;
    • the same harness under compute-sanitizer --tool memcheck with exact-size allocations (I checked that a 1-byte read past the end is flagged): 0 errors, which confirms the "windows never leave the row" claims for the Q8 vector windows and the even-block-count IQ2 record window.
  • Not covered by anything I could run: the __CUDA_ARCH__ >= 900 branches (my GPUs are A100s).

CI is red because of #33994: Windows GPU CUDA CI Pipeline (moe_group_router_impl.cu(157), C4244 → C2220) and the three wasm_* jobs (linear_attention_gates.cc:300,302, -Wshorten-64-to-32). Neither file is touched here; the PR cannot go green until #33994 is fixed.

Non-blocking feedback (inline)

  1. Bound/justify the __CUDA_ARCH__ >= 900 guard and pin the Hopper decoder's arithmetic where CI can run it.
  2. Two NVFP4 tests assert scale bytes >= 0x80, which the documented contract excludes.
  3. GetFetches()[0] after a failed run aborts the whole test binary.

template <bool kAlignedRows>
__device__ __forceinline__ void DecodeNvfp4SubBlock16Bf16(
const uint8_t* row_ptr, int sb32_idx, int half16, uint4& low, uint4& high) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggestion — bound or justify this architecture guard, and pin the Hopper arithmetic where CI can run it. DecodeNvfp4SubBlock16Bf16 (and the staged NVFP4 loop in matmul_packed.cu) take the new path for every __CUDA_ARCH__ >= 900, while the #else comment says other architectures are "not qualified for this subnormal BF16 input encoding". Release builds also target newer architectures (the packaging pipelines build 120-real, and 100-real;103-real;110-real;120-real;121-real on aarch64), which silently get an unqualified path under this condition. Either bound it to what you measured (e.g. >= 900 && < 1000) or state that the encoding is architecture-independent.

I checked the arithmetic independently: DecodeNvfp4WordsBf16 is not arch-guarded, so I compiled it for sm_80 and compared it with the FP32 formula d * (0.5f * kvalues_fp4[code]) rounded to BF16, for every scale byte 0x00–0xFF × (all 16 nibble values at each of the 16 positions + 4096 random 64-bit code words per scale): 0 mismatches in 591,872 in-contract decodes (and 0 in as many out-of-contract ones). So the encoding is exact on SM80 as well, which means the exactness argument can be pinned in CI: the NVFP4 tests run on SM80 but take the FP32 fallback there, so this decoder only executes where a Hopper GPU is available. A small kernel-level test under test/contrib_ops/cuda_kernels/ that launches DecodeNvfp4WordsBf16 over all 128 valid scale bytes × 16 codes × 16 positions against the FP32 formula would run on any SM80+ machine.

Comment on lines +155 to +156
// Signed, non-power-of-two scales plus zero and the 0x7f sentinel.
bytes[group] = static_cast<uint8_t>((col + block * 7 + group * 19) % 255);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggestion (low): this fixture generates scale bytes >= 0x80 ("Signed, non-power-of-two scales"), and the reference below negates d when scale & 128. docs/PackedWeightFormats.md (from #33962) says scale bytes must have bit 7 clear and that bytes >= 0x80 are outside the contract. Nvfp4Bf16PackedOperandExtremes similarly feeds {128, 129, 254, 255} and expects {-0, -2^-9, -448, NaN} (the NaN columns cannot fail anyway: TensorCheck<BFloat16> tests the expected value for NaN, see the earlier #33963 thread).

Both decoders currently agree on the signed-E4M3 reading (in my harness the Hopper decoder reproduces the FP32 decoder's formula bit-for-bit for all 256 scale bytes), so the tests pass, but they pin an unspecified behavior and would fail if the decoder were ever changed to ggml's UE4M3 reading. Either restrict these fixtures to & 0x7f (the Python test already does) or document that bytes >= 0x80 currently decode as signed E4M3 and are not guaranteed.

test.AddInput<float>("weights", {rows, experts}, weights);
test.AddOutput<BFloat16>("output", {rows, hidden}, expected);
RunMoeCuda(test);
const auto* output = test.GetFetches()[0].Get<Tensor>().Data<BFloat16>();

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Robustness: if the preceding RunMoeCuda(test) fails (kernel launch or initialization error), GetFetches() is empty and this [0] aborts the whole test binary instead of failing the test. ORT's Linux builds enable _GLIBCXX_ASSERTIONS, so it is a hard std::vector::operator[] assertion (SIGABRT) that also drops the results of every test that would have run afterwards. I hit exactly this when I forced a CUDA initialization failure (CUDA_FORCE_PTX_JIT=1 on an image without PTX): cublasCreate failed during session init, then __n < this->size() aborted the process inside this test. An ASSERT_EQ(test.GetFetches().size(), 1u) before the access turns it into a normal failure.

@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch 2 times, most recently from 6294c1f to 6d565cf Compare November 10, 2026 04:04
@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch from 6d565cf to c2194e4 Compare November 10, 2026 04:26
Rui Ren (rui-ren) and others added 2 commits November 9, 2026 21:05
…metic

Decode a full sixteen-value NVFP4 group with one shared FP8 scale conversion and two alignment-safe word loads instead of two scalar eight-value decoder calls. Preserve FP32 reconstruction, the complete-operand BF16 cast, MMA operand order, contiguous K partitions and warp reductions. Add independent exact-value production-shape coverage for signed scales, sentinel zero, tails and M1/M7/M8 with initializer and dynamic weights.

Validation: five MatMulPackedPrecisionTest native cases passed without skips through the rebuilt production provider; 524288 exhaustive operand words and complete projection tensors matched bitwise, including offset bindings, padded rows, canaries and M1/M8 row identity. Same-GPU frozen 2048-prompt/512-token ordinary and k7 replays retained all token IDs and workload counters. Large/small CUDA kernels improved from 134.23/56.39 us to 64.49/25.69 us; clang-format and git diff --check passed.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit dbbd73c)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Replace intermediate FP32 operand arrays with byte-permutation BF16 pair packing on SM90+. Qualify word loads only when both weight base and row stride are aligned; retain safe unaligned decoding and the established pre-SM90 fallback. Preserve scale-byte127 zero handling, exact scale/cast operands restored before MMA, activation-dtype dequantization, and unchanged MMA/reduction order. No allocation, initializer PrePack, or dense weight expansion.

Validated the immutable incremental production rebuild: 12 projection/MoE native cases without skips; 524,288 exhaustive operand words and complete-output bitwise equality; M1/M7/M8, offsets, padded rows, tails, signed/extreme scales, sentinel/NaN, and canaries. Same-GPU repeated/reverse-order 2048-prompt fixed-R8 512-output runtime probes preserve tokens/counters with CUDA graphs and zero CUPTI loss. Large projection improves 64.60 to 43.86us, ordinary 69.98 to 75.60 tok/s and k7 417.75 to 442.20. Clang-format and git diff --check pass. Historical 25.45us recovery and production raw-layout PrePack are not claimed.

Signed-off-by: rui-ren <ruiren1225@gmail.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit 7e6379a)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Rui Ren (rui-ren) and others added 8 commits November 9, 2026 21:05
Construct BF16 MMA operand pairs directly from GGML IQ2 records while preserving complete FP32 scale/grid reconstruction and the existing reduction and activation boundaries. Controlled load-form ablations select byte code loads for FC1 and record-word loads for FC2. Add production-width and malformed-route native coverage.

Validation: 7 native tests passed with no skips; complete 17.7M-value bitwise fragment comparisons, full intermediate/final arrays, cohort/offset/canary checks, CUDA memcheck, and clang-format passed. Paired actual-provider replay improved ordinary 54.58 to 57.09 tok/s (+4.59%) and k7 344.82 to 363.82 (+5.51%), with identical full sequences and counters. This is not lossless-k7 release acceptance.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit efdd0c9)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Build an aligned BF16 magnitude codebook from canonical IQ2 byte codes. Every unsigned-byte magnitude is represented exactly; keep full FP32 block-scale/sign reconstruction before the original BF16 operand rounding. Preserve MMA/chunk/reduction order, routing, activation and initialized outputs. Add a positive all-codebook native sweep against the independent canonical CPU table.

Validation: complete bitwise arrays and 17,694,720 fragment values (bindings 0/2, special scales, tails/cohorts and malformed routes); eight native tests, zero skips, and CUDA memcheck with zero errors. Strict C++ formatting/compiler checks pass. Verified actual incremental CUDA provider and complete 512-token sequences/counters in ordinary and k7, CUDA graphs active and no CUPTI losses. Same-GPU v1-v2-v1 bracket improves ordinary 57.124->58.536 tok/s (+2.47%) and k7 363.612->375.139 (+3.17%). No historical-recovery or lossless-release acceptance claimed.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit f9f4013)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Load packed codes, scale bytes and activation words one 16-K half ahead only for the aligned SM90+ KW4/MT2/128-K specialization. Preserve BF16 reconstruction, per-warp MMA order, contiguous K partitions, warp reduction and global-scale-before-cast arithmetic. Retain small-N and unsupported-binding paths after controlled chunk-staging regressions.

Extend the native production-shape fixture with R3 and N-threshold/uneven-K cases. Same-TU production-flag standalone trials show approximately 5% lower R8 N18560/K4096 latency with complete bitwise outputs and exhaustive operand-word equality. The complete production translation unit compiles; native provider and E2E integration remain separate pending qualification gates.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit 23a6934)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Specialize only aligned Q8 contiguous tensor-core projections to read raw 16-byte windows and reassemble the existing 34-byte GGUF records. Retain the existing decoder for offset or stride bindings that do not qualify. Keep FP16 scale reconstruction, BF16 operand rounding, MMA ordering, K partition and epilogue unchanged.

Add focused signed-scale record-window and stride-fallback regression cases. Standalone complete-output comparisons, exhaustive finite-FP16 operand checks, CUDA memcheck and production translation-unit compilation pass; actual provider integration remains a separate qualification gate.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit 8f03dab)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Encode the exact IQ2 magnitude codebook as adjacent BF16 pairs, eliminating byte-to-float magnitude conversions while retaining FP16 scale payloads and complete signed FP32 reconstruction before the single BF16 cast. Preserve existing chunk64 partitions, adjacent-pair MMA ordering, reductions, and epilogues; leave other quantization paths unchanged.

Add independent production-width regression cases with heterogeneous non-power-of-two scales. Qualify full operands and outputs with paired rotating and warm standalone cohorts, compile the production and native-test translation units, and run zero-error full memcheck validation. Actual provider and end-to-end qualification remain separate.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit 20d0faa)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
MatMulPacked accepts the optional global_scale input only for NVFP4. Remove that unsupported input from the new Q8 and IQ2 regression graphs and compare the default unit epilogue instead. Keep their signed non-power-of-two FP16 block scales and complete BF16 operand reconstruction coverage unchanged. The native failures were input-validation errors before kernel execution; no production kernel or API changes are made.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit a0c7e00)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Select aligned activation loads once at launch while retaining the supported offset fallback. Apply IQ2 signs to the complete FP32 product before BF16 conversion without changing scale precision, MMA/K partitions, reductions or squared-ReLU boundaries.

Keep expert-down reconstruction and loading unchanged: applying the same isolated alternatives there regresses distinct-route workloads. Same-TU H100 validation covers complete operands, pipeline intermediates, routing metadata, canaries, offsets and production widths with bitwise identity.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Signed-off-by: rui-ren <ruiren1225@gmail.com>
(cherry picked from commit 391f94e)
Signed-off-by: rui-ren <ruiren1225@gmail.com>
Reuse the dense IQ2 codebook and decoder in MoEPacked while retaining
FC1/FC2 load and sign specializations. Remove the unreachable generic
adjacent packing branch.

Add exact nonuniform MoE oracles and bitwise eight/nine-row equivalence,
a full 256-position IQ2 one-hot sweep, high-nibble NVFP4 coverage on SM80+,
and an internal compiled-MMA-gated 128-K cancellation test. Extend Python
IOBinding coverage for rounded Q8, wide/staged NVFP4, and relu-squared IQ2
misalignment fallbacks.

Preserve lower-layer routing and scale regressions during integration.
Recalibrate non-BF16 route weights so all 36 merged fixture configurations
detect premature BF16 narrowing without weakening the assertion.

Validation: original isolated SM90 projection/MoE and SM80 projection CUDA
compilation, three host objects, 20 ordinary CUDA tests, two Python
IOBinding tests, one internal CUDA cancellation test, and seven CPU
mutation-sensitivity checks passed. Integrated MoE/projection host objects,
36 CPU fixture configurations, configured C++/Python lint and complete
rebase diff checks passed. Integrated GPU execution remains with the parent.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: f4d0e1a3-a389-4cc8-9c85-0fd0e3ffce5b
Signed-off-by: rui-ren <ruiren1225@gmail.com>
@rui-ren
Rui Ren (rui-ren) force-pushed the rui-ren/bolt-packed-latency-on-33994-20261109 branch from c2194e4 to 4f7715c Compare November 10, 2026 05:05

This branch has not been deployed

No deployments
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.

3 participants