[CPU] Add MoEPacked contrib op and AVX2 INT8 GGUF kernels for block-packed weights by BBaditi · Pull Request #33975 · microsoft/onnxruntime

BBaditi · 2026-11-08T07:49:15Z

Description

Adds a native CPU Execution Provider implementation for com.microsoft::MoEPacked along with a header-only 32-wide AVX2 INT8 (_mm256_maddubs_epi16 / vpmaddubsw) GGUF block-quantized dot-product library (gguf_vec_dot_cpu.h) supporting all 12 quantized GGUF formats in gguf_quant_tables.h (Q2_0, Q4_0, Q5_0, Q8_0, Q5_K, Q6_K, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ4_NL, IQ4_XS) plus F32, F16, and BF16.

Key Optimizations

Per-32-Element GgufQ8Block Activation Quantization & 32-Wide AVX2 INT8 (vpmaddubsw):
- Quantizes input activations (x) once per active token and intermediate SwiGLU tiles (inter) on the fly into 32-element GgufQ8Block (int8_t qs[32], float d, int32_t bsum, matching CUDA block_q8_1).
- Evaluates 32 weights per AVX2 instruction using _mm256_maddubs_epi16 + _mm256_madd_epi16 across codebook grid formats (IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS via 2-KB L1D sign LUT + vpsignb), non-linear LUT formats (IQ4_NL, IQ4_XS via pshufb + signed×signed vpsignb), and uniform/superblock formats (Q2_0, Q4_0, Q5_0, Q8_0, Q5_K, Q6_K), deferring horizontal reduction (HsumAvx2) and bsum zero-point correction to the end of each row.
Multi-Row Register Blocking (2x Gate/Up + 4x Down) & 2D Low-Latency Worker Pool:
- VecDotGgufGateUpRow2xQ8Cpu processes 2 intermediate rows (4 weight rows: gate_r0, gate_r1, up_r0, up_r1) in lockstep in Phase 1 (kFc1Tile = 32), and VecDotGgufRow4xQ8Cpu processes 4 output rows in lockstep in Phase 2 (kFc2Tile = 64).
Expert-Grouped Multi-Token Prefill (M > 1) Reuse (ColdExpertGroup):
- Groups multi-token routes by expert_id when M > 1 so all B tokens routed to the same expert reuse each GGUF weight row pair/quad hot in L1D/L2 cache (2.20x–3.60x speedup at M = 16), while keeping a zero-overhead direct path for M = 1 decode.

Single-Layer `Qwen3.8-Flash-Next` (`E=512, H=2560, I=640, k=10`) CPU `MoEPacked` Latency & Parity

Layer (`Gate/Up + Down` Format)	Expert Size (`MB/exp`)	Cosine vs FP32 Dequant Ref	Relative L2 Error	`M=1` CPU `MoEPacked` (`us`)	`M=16` Ungrouped (`ms`)	`M=16` Expert-Grouped (`ms` / Speedup)
`UD-IQ1_S L00` (`IQ2_XXS + Q2_0`)	`1.31 MB`	`0.999935`	`1.14%`	`370.6 us`	`3.92 ms`	`1.09 ms` (`3.60x`)
`UD-IQ1_S L04` (`IQ2_S + Q2_0`)	`1.51 MB`	`0.999929`	`1.19%`	`434.4 us`	`4.18 ms`	`1.42 ms` (`2.94x`)
`UD-IQ2_XS L00` (`IQ2_XXS + IQ4_NL`)	`1.77 MB`	`0.999933`	`1.16%`	`376.3 us`	`4.05 ms`	`1.28 ms` (`3.16x`)
`UD-IQ2_XS L03` (`IQ2_S + IQ4_NL`)	`1.97 MB`	`0.999928`	`1.20%`	`446.6 us`	`4.29 ms`	`1.61 ms` (`2.66x`)
`UD-IQ2_XS L04` (`IQ2_S + Q5_0`)	`2.18 MB`	`0.999930`	`1.18%`	`476.9 us`	`4.41 ms`	`1.78 ms` (`2.48x`)
`UD-IQ3_XXS L14` (`IQ3_XXS + IQ4_NL`)	`2.18 MB`	`0.999934`	`1.15%`	`533.7 us`	`4.30 ms`	`1.95 ms` (`2.20x`)
`UD-IQ3_XXS L47` (`IQ4_XS + IQ4_NL`)	`2.66 MB`	`0.999932`	`1.17%`	`327.6 us`	`3.84 ms`	`1.21 ms` (`3.17x`)

Verification

ctest -R GgufQuantPackedTest: Verifies single-row, fused Gate+Up, 2x Gate/Up, and 4x Down dot products across all 12 quantized GGUF formats (plus F32/F16/BF16) and MoEPacked CPU vs reference dequantization (M = 1 and M = 8 with bit-exact grouped vs ungrouped parity).
pytest onnxruntime/test/python/transformers/test_gguf_packed_ops.py: 11 passed, 174 subtests passed.

Tianlei Wu (tianleiwu) · 2026-11-10T05:21:12Z

Please fix CI failures

Copilot

🟡 Changes recommended

Unchecked bias shapes and scratch-size overflow can cause out-of-bounds memory access, while CUDA silently ignores biases and activation quantization can erase valid floating-point contributions.

10 open issues

Validate optional bias shapes before route execution · New
Avoid Q8 activation quantization for unquantized weights · New
Check scratch-size multiplication for overflow · New
CUDA kernel drops optional biases · New
Parallelize the final weighted combine · New
Prevent overflow in quantizer normalization · New
Preserve NaNs when converting float to BF16 · New
Avoid quadratic token lookup during route mapping · New
Allocate input staging by unique tokens, not routes · New
Add independent reference coverage for the registered operator · New

What changed in this PR

Adds native CPU execution for com.microsoft::MoEPacked, extending ONNX Runtime’s packed-weight mixture-of-experts support beyond CUDA.

Changes:

Implements CPU routing, SwiGLU expert projections, and weighted output combination.
Adds GGUF dot-product helpers with AVX2 INT8 acceleration and expert-grouped prefill execution.
Extends the schema with FP32 activations and optional biases, adds tests, and updates documentation.

File	Description
`onnxruntime/test/contrib_ops/gguf_quant_packed_test.cc`	Adds format-level dot-product checks and CPU grouped/single-token comparisons.
`onnxruntime/core/graph/contrib_ops/contrib_defs.cc`	Adds optional biases and FP32 activation support to the schema.
`onnxruntime/contrib_ops/cpu/quantization/gguf_vec_dot_cpu.h`	Implements quantization, dot products, routing, and grouped expert execution.
`onnxruntime/contrib_ops/cpu/moe/moe_packed_cpu.h`	Declares the templated CPU kernel.
`onnxruntime/contrib_ops/cpu/moe/moe_packed_cpu.cc`	Implements validation, CPU execution, and typed kernel definitions.
`onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc`	Registers the CPU kernels outside minimal builds.
`docs/OperatorKernels.md`	Updates provider kernel listings for MoEPacked.
`docs/ContribOperators.md`	Documents the expanded inputs and activation types.
`.github/workflows/android.yml`	Raises the Android binary-size threshold by 1 KiB.

_{Review effort: Max}

…acked weights Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Copilot

🟡 Changes recommended

The provider test fails to compile in supported MSVC AVX2 builds because it calls a quantizer declared only for GCC/Clang.

10 open issues

10 resolved since last review

_{Review effort: Max}

+#if defined(__AVX2__)
+  GgufQ8Block blk_avx2{};
+  QuantizeBlockQ8Avx2(tiny_x, &blk_avx2);
+  EXPECT_FLOAT_EQ(blk_avx2.d, blk_scalar.d);
+  EXPECT_EQ(blk_avx2.bsum, blk_scalar.bsum);
+  for (int i = 0; i < 32; ++i) {
+    EXPECT_EQ(blk_avx2.qs[i], blk_scalar.qs[i]) << "index=" << i;
+  }
+#endif


+  gguf_quant::ComputeMoEPackedSoftmaxTopKCpu<T>(
+      router_probs->Data<T>(), rows_i, experts_i, k_i, normalize_routing_weights_ != 0,
+      topk_scales.data(), topk_indices.data());


+
+  std::vector<int> route_ids(static_cast<size_t>(total_routes));
+  std::iota(route_ids.begin(), route_ids.end(), 0);
+  std::vector<T> fc2_out(SafeInt<size_t>(total_routes) * hid_i);


+
+  auto* tp = context->GetOperatorThreadPool();
+  auto parallel_for = [tp](std::ptrdiff_t count, const auto& fn) {
+    concurrency::ThreadPool::TryBatchParallelFor(tp, count, fn, 0);


+    const float a = std::fabs(x[i]);
+    if (a > max_abs) max_abs = a;
+  }
+  if (!(max_abs > 0.0f)) {


+#if (defined(__x86_64__) || defined(_M_X64)) && (defined(__GNUC__) || defined(__clang__))
+#define ORT_GGUF_CPU_HAS_AVX2_TARGET 1
+#else
+#define ORT_GGUF_CPU_HAS_AVX2_TARGET 0
+#endif


+      const __m256i ones_16 = _mm256_set1_epi16(1);
+      for (int64_t i = 0; i < nb; ++i) {
+        const uint8_t* blk = weight_row + i * 34;
+        const float s = Fp16BitsToFloatCpu(ReadU16Le(blk)) * x_q8[i].d;


+  for (int u = 0; u < num_unique_toks; ++u) {
+    const T* src_row = input + static_cast<size_t>(unique_toks[u]) * hidden_size;
+    float* dst_row = x_f32 + static_cast<size_t>(u) * hidden_size;
+    for (int j = 0; j < hidden_size; ++j) {
+      dst_row[j] = ActTypeToFloatCpu<T>(src_row[j]);
+    }
+    if (need_x_q8) {
+      QuantizeRowQ8Cpu(dst_row, hidden_size, x_q8 + static_cast<size_t>(u) * q8_blocks_per_hidden);
+    }
+  }


+      std::vector<std::unique_ptr<IExecutionProvider>> cuda_ep;
+      cuda_ep.push_back(DefaultCudaExecutionProvider());
+      tester.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &cuda_ep);


+  RunMoEPackedCpuReferenceTest<MLFloat16>(rng, 1, kGgmlTypeQ8_0, kGgmlTypeQ8_0, kGgmlTypeQ6_K, true, true, 3e-2f);
+  RunMoEPackedCpuReferenceTest<MLFloat16>(rng, 8, kGgmlTypeQ8_0, kGgmlTypeQ8_0, kGgmlTypeQ6_K, false, true, 3e-2f);
+  RunMoEPackedCpuReferenceTest<BFloat16>(rng, 1, kGgmlTypeQ8_0, kGgmlTypeQ8_0, kGgmlTypeQ6_K, true, true, 5e-2f);
+  RunMoEPackedCpuReferenceTest<BFloat16>(rng, 8, kGgmlTypeQ8_0, kGgmlTypeQ8_0, kGgmlTypeQ6_K, false, true, 5e-2f);


need some update

…ed weights (#33976) ### Description Depends on #33975. Extends `cuda::MoEPacked<T>` to implement `IMoeOffloadNode` for hybrid GPU-hot / CPU-cold expert execution on block-packed GGUF MoE models (`Qwen3.8-Flash-Next` `UD-IQ1_S`, `UD-IQ2_XS`, `UD-IQ3_XXS`), integrated with `QMoESessionController` HBM admission planning (`expert-profile.bin` STRP popularity ranking). ### Key Changes - **Slot-Compacted GPU Hot Experts + `expert_to_slot` Indirection**: - Compacts admitted hot experts into contiguous GPU HBM slots while keeping full GGUF expert tensors in zero-copy mmap host memory for CPU cold-expert execution. - Adds an optional `expert_to_slot` map to `LaunchMoEPackedHotExpertsCuda` (`moe_packed.cu`); tokens routed to cold experts (`slot < 0`) are skipped on GPU and accumulated on CPU via `ExecuteMoEPackedColdRoutesCpuWithScratch`. - **Persistent Pinned Staging & True GPU Hot / CPU Cold Overlap (`cudaEvent_t`)**: - Pre-allocates reusable pinned host staging buffers (`pinned_d2h_`, `pinned_h2d_`) and uses `CUDA_EVENT_DISABLE_TIMING` events (`d2h_event_`, `h2d_consumed_event_`) so `LaunchMoEPackedHotExpertsCuda` runs concurrently on the CUDA stream while CPU worker threads execute cold experts. --- ### Table A: 48-Layer `Qwen3.8-Flash-Next` Standard Autoregressive Decode (`M = 1`) | 48-Layer Model (`E=512, k=10`) | Placement (`gpu_experts`) | MoE GPU Weights (`GiB` / `GB`) | MoE VRAM Saved (`GiB` / `GB`) | Hot-Expert Route Hit Rate | Steady-State `M=1` Decode (`ms/tok` / `tok/s`) | Auto-Capped (`~11 GiB` VRAM) `M=1` Decode | Greedy Token Parity (`64/64`) | |---|---|---:|---:|---:|---:|---:|:---:| | **`UD-IQ1_S`** | `100% GPU` (`512`) | `31.10 GiB` (`33.40 GB`) | `—` | `100.0%` | **`20.26 ms` (`49.4 tok/s`)** | `—` | `100%` (Ref) | | **`UD-IQ1_S`** | `75/25 Hybrid` (`384`) | `23.33 GiB` (`25.05 GB`) | **`-7.78 GiB` (`-8.35 GB`)** | `91.0%` | **`21.61 ms` (`46.3 tok/s`)** | `—` | **`64 / 64` (`100%`)** | | **`UD-IQ1_S`** | `50/50 Hybrid` (`256`) | `15.55 GiB` (`16.70 GB`) | **`-15.55 GiB` (`-16.70 GB`)** | `77.3%` | **`23.03 ms` (`43.4 tok/s`)** | **`27.57 ms` (`36.3 tok/s`)** | **`64 / 64` (`100%`)** | | **`UD-IQ2_XS`** | `100% GPU` (`512`) | `39.38 GiB` (`42.28 GB`) | `—` | `100.0%` | **`19.34 ms` (`51.7 tok/s`)** | `—` | `100%` (Ref) | | **`UD-IQ2_XS`** | `75/25 Hybrid` (`384`) | `29.53 GiB` (`31.71 GB`) | **`-9.84 GiB` (`-10.57 GB`)** | `91.6%` | **`20.52 ms` (`48.7 tok/s`)** | `—` | **`64 / 64` (`100%`)** | | **`UD-IQ2_XS`** | `50/50 Hybrid` (`256`) | `19.69 GiB` (`21.14 GB`) | **`-19.69 GiB` (`-21.14 GB`)** | `78.7%` | **`22.00 ms` (`45.5 tok/s`)** | **`29.09 ms` (`34.4 tok/s`)** | **`64 / 64` (`100%`)** | | **`UD-IQ3_XXS`** | `100% GPU` (`512`) | `46.84 GiB` (`50.29 GB`) | `—` | `100.0%` | **`20.55 ms` (`48.7 tok/s`)** | `—` | `100%` (Ref) | | **`UD-IQ3_XXS`** | `75/25 Hybrid` (`384`) | `35.13 GiB` (`37.72 GB`) | **`-11.71 GiB` (`-12.57 GB`)** | `86.8%` | **`21.69 ms` (`46.1 tok/s`)** | `—` | **`64 / 64` (`100%`)** | | **`UD-IQ3_XXS`** | `50/50 Hybrid` (`256`) | `23.42 GiB` (`25.15 GB`) | **`-23.42 GiB` (`-25.15 GB`)** | `68.1%` | **`24.37 ms` (`41.0 tok/s`)** | **`32.38 ms` (`30.9 tok/s`)** | **`64 / 64` (`100%`)** | --- ### Table B: 48-Layer `Qwen3.8-Flash-Next` + `dflash2.onnx` Speculative Decoding (`64` Generated Tokens / Prompt) - **`100% GPU` + `Adaptive M={4,8}`**: **`81.6–94.4 tok/s`** 3-prompt mean (**`1.88x–2.12x`** vs `M=1`), reaching **`140.3–151.3 tok/s` (`2.90x–3.09x`)** on `Adaptive M={4,8}` and **`149.5–157.5 tok/s` (`3.05x–3.23x`)** on `Fixed M=8` for structured sequences. - **`50/50 Hybrid` + `Adaptive M={4,8}`**: **`62.1–74.9 tok/s`** 3-prompt mean (**`1.73x–1.82x`** vs `M=1`), reaching **`88.1–99.1 tok/s` (`2.44x–2.53x`)** on `Adaptive M={4,8}` and **`89.5–103.1 tok/s` (`2.48x–2.63x`)** on `Fixed M=8` — outperforming `100% GPU` `M=1` decoding (`48.7–51.7 tok/s`) while cutting MoE GPU VRAM by `50%`, with **`64/64` (`100%`)** exact token parity. | Model | Placement | Prompt | `M=1` Baseline (`ms/tok` / `tok/s`) | `Fixed M=8` ($\tau$ / `ms/tok` / `tok/s` / `Speedup` / `Match`) | `Fixed M=4` ($\tau$ / `ms/tok` / `tok/s` / `Speedup` / `Match`) | `Adaptive M={4,8}` ($\tau$ / `ms/tok` / `tok/s` / `Speedup` / `Match`) | |---|---|---|---:|---:|---:|---:| | **`UD-IQ1_S`** | `100% GPU` | `P1_structured` | `20.57 ms` (`48.6`) | `7.22` / `6.37 ms` / `156.9` / **`3.23x`** / `100%` | `3.82` / `9.01 ms` / `110.9` / **`2.28x`** / `100%` | `6.50` / `6.67 ms` / `150.0` / **`3.09x`** / `100%` | | **`UD-IQ1_S`** | `100% GPU` | `P2_factual` | `20.47 ms` (`48.9`) | `2.91` / `16.31 ms` / `61.3` / **`1.26x`** / `100%` | `2.37` / `14.49 ms` / `69.0` / **`1.41x`** / `100%` | `2.78` / `13.30 ms` / `75.2` / **`1.54x`** / `100%` | | **`UD-IQ1_S`** | `100% GPU` | `P3_technical` | `20.46 ms` (`48.9`) | `4.00` / `11.88 ms` / `84.2` / **`1.72x`** / `100%` | `2.78` / `12.42 ms` / `80.5` / **`1.65x`** / `100%` | `3.35` / `11.82 ms` / `84.6` / **`1.73x`** / `100%` | | **`UD-IQ1_S`** | `100% GPU` | **`3-Prompt Mean`** | `20.50 ms` (`48.8`) | `4.71` / `11.52 ms` / `86.8` / **`2.07x`** / `100%` | `2.99` / `11.97 ms` / `83.5` / **`1.78x`** / `100%` | `4.21` / `10.59 ms` / `94.4` / **`2.12x`** / `100%` | | **`UD-IQ1_S`** | `50/50 Hybrid` | `P1_structured` | `25.53 ms` (`39.2`) | `7.00` / `9.70 ms` / `103.1` / **`2.63x`** / `100%` | `3.71` / `12.72 ms` / `78.6` / **`2.01x`** / `100%` | `6.30` / `10.09 ms` / `99.1` / **`2.53x`** / `100%` | | **`UD-IQ1_S`** | `50/50 Hybrid` | `P2_factual` | `21.86 ms` (`45.7`) | `2.91` / `18.97 ms` / `52.7` / **`1.15x`** / `100%` | `2.46` / `16.06 ms` / `62.2` / **`1.36x`** / `100%` | `2.78` / `15.30 ms` / `65.3` / **`1.43x`** / `100%` | | **`UD-IQ1_S`** | `50/50 Hybrid` | `P3_technical` | `22.03 ms` (`45.4`) | `3.15` / `16.46 ms` / `60.7` / **`1.34x`** / `100%` | `2.62` / `14.30 ms` / `69.9` / **`1.54x`** / `100%` | `2.86` / `14.64 ms` / `68.3` / **`1.50x`** / `100%` | | **`UD-IQ1_S`** | `50/50 Hybrid` | **`3-Prompt Mean`** | `23.14 ms` (`43.2`) | `4.35` / `15.04 ms` / `66.5` / **`1.71x`** / `100%` | `2.93` / `14.36 ms` / `69.6` / **`1.64x`** / `100%` | `3.98` / `13.35 ms` / `74.9` / **`1.82x`** / `100%` | | **`UD-IQ2_XS`** | `100% GPU` | `P1_structured` | `19.35 ms` (`51.7`) | `7.22` / `6.35 ms` / `157.5` / **`3.05x`** / `100%` | `3.82` / `8.92 ms` / `112.1` / **`2.17x`** / `100%` | `6.50` / `6.61 ms` / `151.3` / **`2.93x`** / `100%` | | **`UD-IQ2_XS`** | `100% GPU` | `P2_factual` | `19.35 ms` (`51.7`) | `2.03` / `23.28 ms` / `43.0` / **`0.83x`** / `100%` | `1.91` / `17.58 ms` / `56.9` / **`1.10x`** / `100%` | `1.91` / `18.52 ms` / `54.0` / **`1.05x`** / `100%` | | **`UD-IQ2_XS`** | `100% GPU` | `P3_technical` | `19.32 ms` (`51.8`) | `3.30` / `14.21 ms` / `70.4` / **`1.36x`** / `100%` | `2.87` / `11.73 ms` / `85.3` / **`1.65x`** / `100%` | `3.30` / `11.63 ms` / `86.0` / **`1.66x`** / `100%` | | **`UD-IQ2_XS`** | `100% GPU` | **`3-Prompt Mean`** | `19.34 ms` (`51.7`) | `4.18` / `14.61 ms` / `68.4` / **`1.75x`** / `100%` | `2.87` / `12.74 ms` / `78.5` / **`1.64x`** / `100%` | `3.90` / `12.25 ms` / `81.6` / **`1.88x`** / `100%` | | **`UD-IQ2_XS`** | `50/50 Hybrid` | `P1_structured` | `25.49 ms` (`39.2`) | `7.22` / `10.11 ms` / `98.9` / **`2.52x`** / `100%` | `3.82` / `12.90 ms` / `77.5` / **`1.98x`** / `100%` | `6.50` / `10.39 ms` / `96.3` / **`2.45x`** / `100%` | | **`UD-IQ2_XS`** | `50/50 Hybrid` | `P2_factual` | `21.46 ms` (`46.6`) | `2.83` / `20.74 ms` / `48.2` / **`1.03x`** / `100%` | `2.41` / `16.84 ms` / `59.4` / **`1.27x`** / `100%` | `2.50` / `17.76 ms` / `56.3` / **`1.21x`** / `100%` | | **`UD-IQ2_XS`** | `50/50 Hybrid` | `P3_technical` | `21.37 ms` (`46.8`) | `3.40` / `15.37 ms` / `65.1` / **`1.39x`** / `100%` | `2.87` / `12.97 ms` / `77.1` / **`1.65x`** / `100%` | `3.40` / `12.92 ms` / `77.4` / **`1.65x`** / `100%` | | **`UD-IQ2_XS`** | `50/50 Hybrid` | **`3-Prompt Mean`** | `22.77 ms` (`43.9`) | `4.48` / `15.41 ms` / `64.9` / **`1.65x`** / `100%` | `3.03` / `14.24 ms` / `70.2` / **`1.63x`** / `100%` | `4.13` / `13.69 ms` / `73.0` / **`1.77x`** / `100%` | | **`UD-IQ3_XXS`** | `100% GPU` | `P1_structured` | `20.64 ms` (`48.5`) | `7.00` / `6.69 ms` / `149.5` / **`3.08x`** / `100%` | `3.71` / `9.61 ms` / `104.0` / **`2.15x`** / `100%` | `6.30` / `7.13 ms` / `140.3` / **`2.90x`** / `100%` | | **`UD-IQ3_XXS`** | `100% GPU` | `P2_factual` | `20.51 ms` (`48.8`) | `4.12` / `11.78 ms` / `84.9` / **`1.74x`** / `100%` | `2.95` / `12.13 ms` / `82.5` / **`1.69x`** / `100%` | `2.95` / `13.50 ms` / `74.1` / **`1.52x`** / `100%` | | **`UD-IQ3_XXS`** | `100% GPU` | `P3_technical` | `20.50 ms` (`48.8`) | `2.86` / `17.09 ms` / `58.5` / **`1.20x`** / `100%` | `2.62` / `13.44 ms` / `74.4` / **`1.53x`** / `100%` | `2.74` / `14.12 ms` / `70.8` / **`1.45x`** / `100%` | | **`UD-IQ3_XXS`** | `100% GPU` | **`3-Prompt Mean`** | `20.55 ms` (`48.7`) | `4.66` / `11.85 ms` / `84.4` / **`2.01x`** / `100%` | `3.10` / `11.73 ms` / `85.3` / **`1.79x`** / `100%` | `4.00` / `11.58 ms` / `86.3` / **`1.96x`** / `100%` | | **`UD-IQ3_XXS`** | `50/50 Hybrid` | `P1_structured` | `27.68 ms` (`36.1`) | `7.00` / `11.17 ms` / `89.5` / **`2.48x`** / `100%` | `3.71` / `14.15 ms` / `70.7` / **`1.96x`** / `100%` | `6.30` / `11.35 ms` / `88.1` / **`2.44x`** / `100%` | | **`UD-IQ3_XXS`** | `50/50 Hybrid` | `P2_factual` | `28.17 ms` (`35.5`) | `4.12` / `19.80 ms` / `50.5` / **`1.42x`** / `100%` | `2.95` / `18.87 ms` / `53.0` / **`1.49x`** / `100%` | `2.95` / `21.55 ms` / `46.4` / **`1.31x`** / `100%` | | **`UD-IQ3_XXS`** | `50/50 Hybrid` | `P3_technical` | `22.27 ms` (`44.9`) | `3.05` / `17.84 ms` / `56.1` / **`1.25x`** / `100%` | `2.62` / `14.91 ms` / `67.1` / **`1.49x`** / `100%` | `2.74` / `15.43 ms` / `64.8` / **`1.44x`** / `100%` | | **`UD-IQ3_XXS`** | `50/50 Hybrid` | **`3-Prompt Mean`** | `26.04 ms` (`38.4`) | `4.72` / `16.27 ms` / `61.5` / **`1.72x`** / `100%` | `3.10` / `15.97 ms` / `62.6` / **`1.65x`** / `100%` | `4.00` / `16.11 ms` / `62.1` / **`1.73x`** / `100%` | --- ### Notes on Multi-Token (`M > 1`) Speculative Verification 1. **Double-Buffered GPU State Banks (`banks[0]` / `banks[1]`) + `libqwen38_plane.so` (`0.19 ms` Single-Launch Replay)**: - In `CausalConvWithState` (`causal_conv_with_state_impl.cu`), when `seq_len > 1` and `kernel_size = 10` (`past_key_values.1.ple_conv_state`), `CausalConvProjectionKernel` launches with `gridDim.z = seq_len` and `blockIdx.z == 0` writes `present` in place while `blockIdx.z > 0` reads `past`. Binding distinct double-buffered `past` (`banks[cur_bank]`) and `present` (`banks[1 - cur_bank]`) state banks eliminates this in-place aliasing race and removes pre-verify D2D state copy overhead, while `qwen38_plane_replay_state_updates` replays the accepted prefix (`kept_count = k_acc`) across all 36 `GatedDeltaNet` layers, 37 `CausalConvWithState` states, and 12 `PackedSparseAttentionIndexer` buffers in a single `0.19 ms` launch. 2. **SM80 (`A100`) `RouterGateUpProj` `cublasGemmEx` Dispatch**: - On SM80 (`A100`), `RouterGateUpProj` dispatches `cublasGemmEx` (`CUBLAS_COMPUTE_32F`), which selects a SIMT GEMV kernel at `rows = 1` vs an `mma.sync.m16n8k16` Tensor Core GEMM kernel at `rows >= 2`. Evaluating all speculative horizons (`Fixed M=8`, `Fixed M=4`, `Adaptive M={4,8}`) against the `kept_count = 1` Tensor Core greedy reference (matching `ORT_QWEN38_FORCE_REJECT=1` in `onnxruntime-genai`) yields **`64/64` (`100%`)** exact greedy token sequence parity across all models, placements, and prompts. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…and GPU FC2 decode (#34067) > **Note**: This PR is stacked on #33975 (`qmoe-packed-gguf-cpu-ep`) and #33976 (`qmoe-packed-gguf-cpu-offload`), and pairs with #34066 (`qmoe-packed-gguf-cuda-matmul-m8-opt` on `main`). ## Summary This PR optimizes `com.microsoft::MoEPacked` multi-token (`M = 4, 8`) speculative decoding verification across both the CPU hybrid cold-route path (`gguf_vec_dot_cpu.h`) and the GPU Tensor Core FC2 decode path (`moe_packed.cu`): 1. **CPU Multi-Token Cold-Route Execution (`ExecuteMoEPackedColdRoutesCpuWithScratch` in `gguf_vec_dot_cpu.h`)**: - **Zero-Heap Stack Buffers**: Replaces per-layer `std::vector<int> sorted_c` and `std::vector<ColdExpertGroup> groups` allocations with stack buffers (`sorted_c_stack[128]`, `groups_stack[128]`) for `num_cold_routes <= 128`, eliminating `96` heap allocations/frees per 48-layer verify step. - **Parallel Multi-Token Input Q8 Quantization**: Parallelizes `BF16 -> FP32 + QuantizeRowQ8Cpu` across `num_unique_toks` when `num_unique_toks >= 2`. - **2-Route Fused AVX2 Grouped Dot Products (`VecDotGgufRow4xQ8Batch2Avx2` & `VecDotGgufGateUpRow2xQ8Batch2Cpu`)**: When $\ge 2$ tokens in an `M = 4` or `M = 8` verification batch route to the same cold expert (`ColdExpertGroup` with `grp.count >= 2`, where measured cold-route reuse $C/U$ is `1.39x-1.48x` at `M=4` and `1.86x-2.07x` at `M=8`), each GGUF weight sub-block (`IQ2_XXS`, `IQ2_XS`, `IQ2_S`, `IQ3_XXS`, `IQ4_NL`, `IQ4_XS`, `Q2_0`) is unpacked once in AVX2 registers and multiplied against two routes' Q8 activations simultaneously with `0` ULP difference. 2. **GPU Tensor Core FC2 Decode/Verify (`MoEPackedFc2Bf16TcKernel<1>` in `moe_packed.cu`)**: - Eliminates the `10.75 KB` `s_act_u4` shared-memory activation staging buffer, integer-division staging loop, and K-loop `__syncthreads()` barrier when `kMaxSubTiles == 1` (`total_routes <= 128`), loading `fc1_out` directly via 128-bit `__ldg`. - Stages GGUF lookup tables (`iq2xxs_grid`, `iq2xs_grid`, `iq2s_grid`, `iq3xxs_grid`) in shared memory (`s_lut`, `DecodeGgufSubBlock8SmemLutTpl<kDownType>`) and switches decode CTA geometry from `8` warps (`128` output channels/CTA) to `4` warps (`64` output channels/CTA) in both `DispatchMoEPackedHotExpertsCuda` and `LaunchMoEPackedFc2Bf16TcDispatch<1>`, doubling CTA grid parallelism across SMs. --- ## Measured Performance Impact (48-Layer `Qwen3.8-Flash-Next` + `dflash2.onnx` on NVIDIA A100-SXM4-80GB) ### Table A: 48-Layer Verify Step Latency Progression (`UD-IQ2_XS`) | Metric (`UD-IQ2_XS`) | Stage 0 (Baseline) | Stage 1 (This PR: `MoEPacked` FC2 + CPU Batch-2) | Stage 2 (+ PR #34066 `MatMulPacked` & `ReadU32Dev`) | Total Gain | |---|---:|---:|---:|---:| | **`100% GPU` `MoEPacked` (48 nodes, `M=1 / M=4 / M=8`)** | `10.74 / 15.82 / 20.29 ms` | **`10.23 / 15.05 / 18.81 ms`** | **`10.28 / 15.32 / 19.24 ms`** | **`-0.5 to -1.5 ms`** | | **`50/50 Hybrid` `MoEPacked` (48 nodes, `M=1 / M=4 / M=8`)** | `17.40 / 27.02 / 37.81 ms` | **`18.31 / 24.69 / 33.91 ms`** | **`16.98 / 24.83 / 34.33 ms`** | **`-2.3 to -3.9 ms`** | | **`100% GPU` `T_verify(M=4)`** | `34.48-35.10 ms` | `29.02-29.73 ms` | **`26.46-27.32 ms`** | **`-7.9 ms/step` (`-22.8%`)** | | **`100% GPU` `T_verify(M=8)`** | `50.32-51.68 ms` | `35.76-38.05 ms` | **`32.90-35.27 ms`** | **`-16.9 ms/step` (`-33.2%`)** | | **`50/50 Hybrid` `T_verify(M=4)`** | `38.68-52.97 ms` | `31.73-41.84 ms` | **`29.64-39.80 ms`** | **`-11.1 ms/step` (`-24.2%`)** | | **`50/50 Hybrid` `T_verify(M=8)`** | `56.08-80.69 ms` | `41.46-56.03 ms` | **`38.69-53.59 ms`** | **`-22.2 ms/step` (`-32.5%`)** | ### Table B: End-to-End `dflash2.onnx` Speculative Decoding (`64` Generated Tokens / Prompt, Combined Stage 2) | Model | Placement | `M=1` Decode (`ms/tok` / `tok/s`) | `Fixed M=8` Peak (`P1_structured`) | `Adaptive M={4,8}` 3-Prompt Mean (`tok/s` / `Speedup`) | Greedy Token Parity | |---|---|---:|---:|---:|---:| | **`UD-IQ1_S`** | `100% GPU` | `21.16 ms` (`47.3 tok/s`) | **`185.3 tok/s` (`3.93x`)** | **`106.0 tok/s` (`86.2` harm., `2.24x`)** | **`64/64` (`100.0%`)** | | **`UD-IQ1_S`** | `50/50 Hybrid` (`-5.21 GB`) | `24.27 ms` (`41.2 tok/s`) | **`129.6 tok/s` (`3.43x`)** | **`90.2 tok/s` (`84.9` harm., `2.23x`)** | **`64/64` (`100.0%`)** | | **`UD-IQ2_XS`** | `100% GPU` | `19.91 ms` (`50.2 tok/s`) | **`197.4 tok/s` (`3.93x`)** | **`123.7 tok/s` (`111.3` harm., `2.46x`)** | **`64/64` (`100.0%`)** | | **`UD-IQ2_XS`** | `50/50 Hybrid` (`-5.78 GB`) | `23.77 ms` (`42.1 tok/s`) | **`126.8 tok/s` (`3.32x`)** | **`89.8 tok/s` (`82.7` harm., `2.17x`)** | **`64/64` (`100.0%`)** | | **`UD-IQ3_XXS`** | `100% GPU` | `22.27 ms` (`44.9 tok/s`) | **`177.5 tok/s` (`3.96x`)** | **`109.3 tok/s` (`98.1` harm., `2.44x`)** | **`64/64` (`100.0%`)** | | **`UD-IQ3_XXS`** | `50/50 Hybrid` (`-7.04 GB`) | `27.77 ms` (`36.0 tok/s`) | **`115.7 tok/s` (`3.30x`)** | **`80.1 tok/s` (`74.1` harm., `2.22x`)** | **`64/64` (`100.0%`)** | --- ## Validation - `/usr/bin/clang-format --dry-run --Werror` and `git diff --check`: **0 issues** - `ctest --output-on-failure -V -R "GgufQuantPackedTest"`: **2/2 passed (`100%`)** - `pytest onnxruntime/test/python/transformers/test_gguf_packed_ops.py -v`: **33 passed, 811 subtests passed (`100%`)** Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…16C/Batch4/40-core cold decode (#34114) ## Description Stacks on #33976 (and integrates #34066 and #34067). Optimizes `com.microsoft::MoEPacked` CPU/GPU hybrid expert offload and `com.microsoft::MatMulPacked` multi-token (`M=2..8`) MTP verification for block-packed GGUF MoE models (`Qwen3.8-Flash-Next` `UD-IQ3_XXS`) based on Nsight Systems (`nsys --cuda-graph-trace=node`) timeline profiling on NVIDIA H200: 1. **Prefill**: Eliminates cross-stream H2D Copy-Engine head-of-line blocking by double-buffering per-wave route/tile metadata (`d_wave_sorted_routes[2]`, `d_wave_tiles[2]`, `d_wave_num_tiles[2]`) on `wave_staging.copy_stream` and decoupling host `std::memcpy` (`h2d_done[b]`) from device compute (`compute_done[b]`), increasing true 2-stream compute/H2D overlap **3.77x** (`6.6% -> 24.9%`, `154.4 ms -> 478.8 ms` at `8K`) and cutting `2K–128K` TTFT by **`-14.7% to -25.9%`** (`+17.3% to +35.0%` prefill tok/s, saving `8.67s–9.72s` at `128K`). 2. **Decode**: Replaces 129 software `MLFloat16::ToFloat()` scale conversions in `gguf_vec_dot_cpu.h` AVX2 loops with hardware F16C `_mm_cvtph_ps`, adds 4-route (`Batch-4`) fused AVX2 INT8 (`vpmaddubsw`) dot-product kernels (`AccumIq2S32Q8Batch4Avx2`, `AccumIq3Xxs32Q8Batch4Avx2`, `AccumIq4Nl32Q8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Avx2`, `VecDotGgufGateUpRow2xQ8Batch4Cpu`), scales `MoEPackedParallelPool` to `40` physical cores, and dispatches `MatMulPackedBf16TcWarpKSplit<4, 1>` for `N <= 2560`, cutting `MoEPackedDoorbellWaitKernel` by **`-49.6%`** (`53.40 -> 26.91 ms/cycle` exact; **`-81.9%`** to `9.66 ms/cycle` with `tail_sub=0.02`) and accelerating decode throughput by **`+60.1% to +79.4%`** (exact routing) and **`+102.8% to +152.5%` (`2.03x–2.53x`)** with tail expert substitution. ## Summary of Changes ### 1. `MoEPacked` 3-Stage Prefill Pipeline & Cold-Expert Decode (`onnxruntime/contrib_ops/cuda/moe/moe_packed.cc`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cc` | Double-buffer `d_wave_sorted_routes[2]`, `d_wave_tiles[2]`, and `d_wave_num_tiles[2]` in `MoEPackedSharedPrefillWaveStaging` and copy them on `wave_staging.copy_stream` prior to recording `h2d_done[b]`; synchronize host `std::memcpy` only on `h2d_done[b]` while `copy_stream` waits on `compute_done[b]` on the device timeline; parallelize host expert staging across `w_cnt * 3` tasks; raise `MoEPackedParallelPool::ResolveThreads()` default cap from `24` to `40` cores; integrate `Batch-4` + `Batch-2` Q8 cold-expert execution in `RunColdDecodeStep()` (opt-in via `ep.cuda.moe_packed_cpu_q8_dot=1` or `ORT_MOE_PACKED_CPU_Q8_DOT=1` so existing FP16/BF16 models retain full FP32 cold-expert dot precision by default). | ### 2. Hardware F16C & 4-Route (`Batch-4`) AVX2 INT8 GGUF Kernels (`onnxruntime/contrib_ops/cpu/quantization/gguf_vec_dot_cpu.h`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cpu/quantization/gguf_vec_dot_cpu.h` | Add `Fp16ToFloatF16c` (`_mm_cvtph_ps`) under `ORT_GGUF_CPU_HAS_AVX2_TARGET` and replace 129 software `Fp16BitsToFloatCpu` calls inside AVX2 loops; add 4-route fused AVX2 INT8 dot-product kernels (`AccumIq2S32Q8Batch4Avx2`, `AccumIq3Xxs32Q8Batch4Avx2`, `AccumIq4Nl32Q8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Avx2`, `VecDotGgufRow4xQ8Batch4Cpu`, `VecDotGgufGateUpRow2xQ8Batch4Cpu`) plus 2-route `Batch-2` kernels from #34067. | ### 3. `MoEPacked` Decode FC2 & `MatMulPacked` Multi-Token GEMV (`moe_packed.cu`, `matmul_packed.cu`, `gguf_quant_cuda.cuh`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_packed.cu` | Integrate #34067 4-warp (`64` output channels/CTA) `MoEPackedFc2Bf16TcKernel<1>` decode geometry with shared-memory `s_lut` GGUF lookup tables and direct `__ldg(fc1_row)` activation loads. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.cu` | Integrate #34066 `(NR=2, KW=2)` multi-token GEMV and dispatch `MatMulPackedBf16TcWarpKSplit<4, 1>` when `N <= 2560` (`M in 2..8`) to double CTA grid occupancy (`1,280 -> 2,560` CTAs across 132 H200 SMs, `-45.4%` per-call latency on 72 GDN projection calls/cycle). | | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh` | Integrate #34066 uniform `uint16x2` `__ldg` loads in `ReadU32Dev` to eliminate 50% warp divergence on 2-byte-aligned GGUF blocks (`Q8_0`, `Q6_K`, `IQ4_NL`, `IQ4_XS`). | --- ## Performance & Accuracy Evaluation (`Qwen3.8-Flash-Next` `UD-IQ3_XXS` + `dflash2` `S=8` MTP on H200) ### 1. Setting 1: `30 GiB VRAM + 50 GiB RAM` (`device_expert_bytes = 12.50 GiB`, `6,463` GPU Resident Slots, Peak VRAM = `29.82 GiB`, Peak `VmRSS` = `35.44–36.09 GiB`) | Configuration | `2K` TTFT / Decode | `8K` TTFT / Decode | `32K` TTFT / Decode | `128K` TTFT / Decode | MMLU-Pro Smoke (`20`) Acc | MMLU Tok-Wtd / Med TPS | CPU `ms/step` (`us/exp`) | |---|---|---|---|---|---:|---|---:| | **1. Baseline (`origin/main` `6cc3c09506`)** | `1.061s` (`1,931 t/s`) / `39.18 t/s` | `2.433s` (`3,367 t/s`) / `38.87 t/s` | `9.156s` (`3,579 t/s`) / `39.28 t/s` | `38.790s` (`3,379 t/s`) / `38.31 t/s` | `100.0%` (`20/20`) | `41.36` / `43.84 t/s` | `61.58 ms` (`43.72 us`) | | **2. Merged PRs (#33975+#33976+#34066+#34067, FP32 CPU)** | `1.052s` (`1,947 t/s`) / `43.48 t/s` | `2.432s` (`3,369 t/s`) / `44.26 t/s` | `8.961s` (`3,657 t/s`) / `43.09 t/s` | `38.357s` (`3,417 t/s`) / `38.55 t/s` | `100.0%` (`20/20`) | `42.57` / `48.49 t/s` | `61.65 ms` (`43.52 us`) | | **3. Merged PRs (#33975+#33976+#34066+#34067, Q8 Batch-2 CPU)** | `1.035s` (`1,978 t/s`) / `50.17 t/s` | `2.395s` (`3,420 t/s`) / `50.37 t/s` | `8.942s` (`3,665 t/s`) / `49.13 t/s` | `38.318s` (`3,421 t/s`) / `47.88 t/s` | `100.0%` (`20/20`) | `49.50` / `57.69 t/s` | `47.85 ms` (`33.35 us`) | | **4. This PR (Exact Routing)** | **`0.786s` (`2,607 t/s`, `+35.0%`) / `62.73 t/s` (`+60.1%`)** | **`1.807s` (`4,534 t/s`, `+34.6%`) / `69.30 t/s` (`+78.3%`)** | **`6.813s` (`4,809 t/s`, `+34.4%`) / `69.45 t/s` (`+76.8%`)** | **`30.124s` (`4,351 t/s`, `+28.8%`) / `61.87 t/s` (`+61.5%`)** | `95.0%` (`19/20`) | **`67.12` / `79.32 t/s` (`+80.9%`)** | **`32.80 ms` (`22.89 us`)** | | **5. This PR + TailSub (`thr=0.02, min_rank=2`)** | **`0.853s` (`2,400 t/s`, `+24.3%`) / `97.11 t/s` (`+147.9%`)** | **`1.969s` (`4,160 t/s`, `+23.5%`) / `88.70 t/s` (`+128.2%`)** | **`7.495s` (`4,372 t/s`, `+22.2%`) / `81.71 t/s` (`+108.0%`)** | **`32.830s` (`3,993 t/s`, `+18.2%`) / `77.68 t/s` (`+102.8%`)** | **`100.0%` (`20/20`)** | **`97.76` / `108.06 t/s` (`+146.5%`)** | **`14.13 ms` (`25.95 us`)** | ### 2. Setting 2: `22 GiB VRAM + 24 GiB RAM` (`device_expert_bytes = 5.60 GiB`, `2,843` GPU Resident Slots, Peak VRAM = `21.71 GiB`, Peak `RssAnon` = `0.91–0.95 GiB`) | Configuration | `2K` TTFT / Decode | `8K` TTFT / Decode | `32K` TTFT / Decode | `128K` TTFT / Decode | MMLU-Pro Smoke (`20`) Acc | MMLU Tok-Wtd / Med TPS | CPU `ms/step` (`us/exp`) | |---|---|---|---|---|---:|---|---:| | **1. Baseline (`origin/main` `6cc3c09506`)** | `1.216s` (`1,684 t/s`) / `29.39 t/s` | `4.125s` (`1,986 t/s`) / `31.86 t/s` | `15.739s` (`2,082 t/s`) / `31.76 t/s` | `66.065s` (`1,984 t/s`) / `30.32 t/s` | `95.0%` (`19/20`) | `34.00` / `39.65 t/s` | `74.14 ms` (`42.04 us`) | | **2. Merged PRs (#33975+#33976+#34066+#34067, FP32 CPU)** | `1.197s` (`1,712 t/s`) / `34.54 t/s` | `4.042s` (`2,027 t/s`) / `35.58 t/s` | `15.762s` (`2,079 t/s`) / `33.70 t/s` | `66.006s` (`1,986 t/s`) / `34.26 t/s` | `95.0%` (`19/20`) | `37.03` / `42.90 t/s` | `73.66 ms` (`42.24 us`) | | **3. Merged PRs (#33975+#33976+#34066+#34067, Q8 Batch-2 CPU)** | `1.230s` (`1,665 t/s`) / `37.75 t/s` | `4.028s` (`2,034 t/s`) / `41.31 t/s` | `15.600s` (`2,101 t/s`) / `36.93 t/s` | `66.186s` (`1,980 t/s`) / `41.27 t/s` | `95.0%` (`19/20`) | `43.87` / `49.14 t/s` | `57.36 ms` (`32.47 us`) | | **4. This PR (Exact Routing)** | **`0.980s` (`2,090 t/s`, `+24.1%`) / `52.61 t/s` (`+79.0%`)** | **`3.345s` (`2,449 t/s`, `+23.3%`) / `52.73 t/s` (`+65.5%`)** | **`13.205s` (`2,482 t/s`, `+19.2%`) / `56.98 t/s` (`+79.4%`)** | **`56.343s` (`2,326 t/s`, `+17.3%`) / `50.16 t/s` (`+65.4%`)** | `95.0%` (`19/20`) | **`62.68` / `70.61 t/s` (`+78.1%`)** | **`38.21 ms` (`21.71 us`)** | | **5. This PR + TailSub (`thr=0.02, min_rank=2`)** | **`0.962s` (`2,130 t/s`, `+26.5%`) / `74.21 t/s` (`+152.5%`)** | **`3.321s` (`2,467 t/s`, `+24.2%`) / `74.47 t/s` (`+133.7%`)** | **`13.191s` (`2,484 t/s`, `+19.3%`) / `68.41 t/s` (`+115.4%`)** | **`56.890s` (`2,304 t/s`, `+16.1%`) / `72.87 t/s` (`+140.3%`)** | `95.0%` (`19/20`) | **`89.00` / `100.57 t/s` (`+153.6%`)** | **`17.55 ms` (`23.32 us`)** | ### 3. Nsight Systems (`nsys --cuda-graph-trace=node`) Before-vs-After Verification (`30g_50g`, `8192` Context) | Nsight Systems Metric (`8192` Context, `17` `S=8` MTP Cycles) | Pre-Optimization | Post-Opt (Exact) | Post-Opt (`tail-sub=0.02`) | Delta (Post Exact / Tail-Sub vs. Pre) | |---|---:|---:|---:|---| | **8K Prefill GPU Window (`Chunk 0 + Chunk 1`)** | `2,354.70 ms` | **`1,922.92 ms`** | **`1,919.95 ms`** | **`-431.78 ms` (`-18.3%`) / `-434.75 ms` (`-18.5%`)** | | **True 2-Stream Compute/H2D Overlap (`stream` & `copy_stream` both active)** | `154.43 ms` (`6.6%`) | **`478.83 ms` (`24.9%`)** | **`480.18 ms` (`25.0%`)** | **`+324.40 ms` (`3.10x` time / `3.77x` share)** | | **Both GPU Streams Idle in Prefill** | `293.64 ms` (`12.5%`) | **`84.03 ms` (`4.4%`)** | **`83.47 ms` (`4.3%`)** | **`-209.61 ms` (`-71.4%` idle bubble)** | | **48-Layer Target Verify CUDA Graph Duration (`ms/pass`)** | `84.11 ms/pass` | **`53.36 ms/pass`** | **`38.91 ms/pass`** | **`-30.75 ms` (`-36.6%`) / `-45.20 ms` (`-53.7%`)** | | **`MoEPackedDoorbellWaitKernel` (`48` calls/cycle)** | `53.40 ms/cyc` (`60.6%`) | **`26.91 ms/cyc` (`47.2%`)** | **`9.66 ms/cyc` (`22.7%`)** | **`-49.6%` (`1.98x` shorter) / `-81.9%` (`5.53x` shorter)** | | **`MatMulPackedBf16TcWarpKSplit<4, 1>` (`72` GDN calls/cycle, `N <= 2560`)** | `9.44 ms/cyc` (`131.1 us/call` on `<4,2>`) | **`5.15 ms/cyc` (`71.6 us/call`)** | **`5.17 ms/cyc` (`71.8 us/call`)** | **`-4.29 ms/cyc` (`-45.4%` per-call speedup)** | ## Testing - `pytest onnxruntime/test/python/transformers/test_gguf_packed_ops.py`: passes CPU/CUDA/hybrid offload numerical parity and tail-substitution tests on H200 (`sm_90`). - End-to-end `2K / 8K / 32K / 128K` context benchmarks and 20-item MMLU-Pro smoke evaluation under both `30 GiB VRAM + 50 GiB RAM` and `22 GiB VRAM + 24 GiB RAM` settings. ## Checklist - [x] Tests verified (`test_gguf_packed_ops.py` and MMLU-Pro smoke evaluation) - [x] No breaking changes - [x] Documentation updated --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

BBaditi mentioned this pull request Nov 8, 2026

[CUDA/CPU] Add hybrid GPU/CPU expert offload for MoEPacked block-packed weights #33976

Merged

github-advanced-security Bot found potential problems Nov 8, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cpu/moe/moe_packed_cpu.cc Fixed

BBaditi force-pushed the qmoe-packed-gguf-cpu-ep branch 2 times, most recently from d3283ae to 5ac6332 Compare November 9, 2026 06:17

Base automatically changed from tlwu/qwen38-gguf to main November 9, 2026 17:56

BBaditi force-pushed the qmoe-packed-gguf-cpu-ep branch from 5ac6332 to 85653e0 Compare November 10, 2026 02:20

BBaditi mentioned this pull request Nov 10, 2026

[MoE] Optimize multi-token (M=4, 8) MoEPacked CPU hybrid cold routes and GPU FC2 decode #34067

Merged

Tianlei Wu (tianleiwu) added the qmoe_offload label Nov 10, 2026

BBaditi force-pushed the qmoe-packed-gguf-cpu-ep branch 2 times, most recently from 3cf8acc to feddfe1 Compare November 10, 2026 18:26

Copilot started work on behalf of BBaditi November 10, 2026 20:49

View session

Copilot finished work on behalf of BBaditi November 10, 2026 20:53

Tianlei Wu (tianleiwu) mentioned this pull request Nov 11, 2026

[CUDA] Optimize MoEPacked GGUF offload 3-stage prefill pipeline and F16C/Batch4/40-core cold decode #34114

Merged

3 tasks done

Tianlei Wu (tianleiwu) requested a max review from Copilot November 11, 2026 21:48

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) November 11, 2026 21:49

View session

Copilot AI reviewed Nov 11, 2026

View reviewed changes

[CPU] Add MoEPacked contrib op and AVX2 INT8 GGUF kernels for block-p…

8ec9abb

…acked weights Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

BBaditi force-pushed the qmoe-packed-gguf-cpu-ep branch from a8a946b to 8ec9abb Compare November 12, 2026 01:07

Tianlei Wu (tianleiwu) requested a max review from Copilot November 12, 2026 02:28

Tianlei Wu (tianleiwu) marked this pull request as ready for review November 12, 2026 02:28

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) November 12, 2026 02:28

View session

Copilot AI reviewed Nov 12, 2026

View reviewed changes

Tianlei Wu (tianleiwu) enabled auto-merge (squash) November 12, 2026 03:09

Tianlei Wu (tianleiwu) previously approved these changes Nov 12, 2026

View reviewed changes

Navigation Menu

Repository navigation

[CPU] Add MoEPacked contrib op and AVX2 INT8 GGUF kernels for block-packed weights - #33975

[CPU] Add MoEPacked contrib op and AVX2 INT8 GGUF kernels for block-packed weights#33975
BBaditi wants to merge 1 commit into
mainfrom
qmoe-packed-gguf-cpu-ep

BBaditi commented Nov 8, 2026 •

edited

Loading

Uh oh!

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 10, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

BBaditi commented Nov 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Key Optimizations

Single-Layer Qwen3.8-Flash-Next (E=512, H=2560, I=640, k=10) CPU MoEPacked Latency & Parity

Verification

Uh oh!

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 10, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

BBaditi commented Nov 8, 2026 •

edited

Loading

Single-Layer `Qwen3.8-Flash-Next` (`E=512, H=2560, I=640, k=10`) CPU `MoEPacked` Latency & Parity