[BBaditi]

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.