Repository navigation
[CUDA] Optimize QMoE CPU/GPU offload for Windows WDDM and AVX2 INT2/INT3 cold experts - #33916
Conversation
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved issues can falsely exhaust pinned memory, exceed explicit device-memory budgets, and override per-run warmup controls or terminate the process on malformed input, while the newly enabled SM89 path lacks parity and dispatch test coverage.
4 open issues
What changed in this PR
This PR improves ONNX Runtime’s mixed-precision QMoE CPU/GPU offload, targeting Windows WDDM memory pressure and prefill/decode performance on consumer GPUs.
Changes:
- Updates file-backed expert loading, session memory planning, and pinned-memory reclamation.
- Expands streaming tiles and double-buffered staging, with Windows prefetch enabled by default.
- Optimizes AVX2 INT2/INT3 expert computation, enables SM89 INT8 Tensor-Core dispatch, and resolves MSVC variable-shadowing warnings.
| File | Description |
|---|---|
onnxruntime/core/providers/cuda/qmoe_session_controller.h |
Adds machine-aware budgets and Windows mapping policies. |
onnxruntime/contrib_ops/cuda/moe/qmoe_offload.h |
Stores expert-bank file paths and offsets. |
onnxruntime/contrib_ops/cuda/moe/qmoe_offload.cc |
Updates cache initialization, pinned budgets, streaming, prefetch, and warmup handling. |
onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h |
Enables INT8 Tensor-Core dispatch on SM89 and newer devices. |
onnxruntime/contrib_ops/cuda/bert/paged_attention_impl.cu |
Renames local constants to avoid MSVC shadowing warnings. |
onnxruntime/contrib_ops/cpu/moe/qmoe_shared_layout.cc |
Accepts additional Windows mapped expert-bank sources. |
onnxruntime/contrib_ops/cpu/moe/qmoe_shared_layout_int3_avx2.cc |
Adds eight-column computation and reuses activation loads. |
onnxruntime/contrib_ops/cpu/moe/qmoe_shared_layout_int2_avx2.cc |
Adds an eight-column prepacked AVX2 path with weight scaling before products. |
Review effort: Balanced
…NT3 cold experts 1. Fix Windows FileExpertStore::InitializeMapped and QMoEOffload::InitializeBoundedHostCache: - Propagate external_file_path and external_file_offset from MoeExpertSourceDescriptor into QMoEOffload::InitializeBoundedHostCache. - In FileExpertStore::InitializeMapped on Windows (_WIN32), accept committed MEM_MAPPED (or read_opt_enabled) views via VirtualQuery when file_path is null. 2. Fix pinned memory budget accounting and persistent decode pinned starvation: - Add SessionPromotionBudget::HeldBytes() and compute reserved = bounce_held + max(decode_held, reserve_target) in QMoEOffload::ComputeShared so promotion bounce buffers are not conflated with persistent decode headroom. - Release idle promotion bounce and landing-ring pinned staging in QMoEOffload::EnsurePersistentDecodeState when persistent decode pinned reservation needs headroom. 3. Automatic machine-aware session planning on Windows WDDM & RAM-constrained hosts (QMoESessionController::PlanAndAdmitSession): - Automatically prefer direct OS file-backed mapping (use_bounded_file_cache = false) on Windows when ullTotalPhys < host_expert_bytes + 24 GiB, avoiding multi-GiB private heap commits and enabling 2-bank parallel pinned wave staging. - Automatically size implicit device_working_bytes (560 MiB) and pinned_bytes (160 MiB) on full-model sessions (num_layers >= 24, total_vram >= 16 GiB, total_full_gpu_bytes >= 8 GiB), and adapt admitted_gpu_expert_budget against platform_free_vram and SessionBudgetAccount without spilling into WDDM non-local memory. - Support Windows MEM_MAPPED views in EvictUploadedHostPages via VirtualQuery + VirtualUnlock. 4. Enable 8K single-tile streaming prefill, 2-bank double-buffered wave staging, and default Windows PrefetchVirtualMemory across INT2/INT3/INT4 layers: - Default stream_rows to 8192 for INT2, INT3, and full-budget INT4 layers. - Automatically halve wave_experts (32 -> 16) in ComputeShared when needed to fit 2 weight staging banks (stage_weight_banks = 2). - Extend ORT_QMOE_WIN_PREFETCH_MMAP (default 1 on Windows) to all file-mapped integer-bank layers. 5. CPU AVX2 INT2/INT3 cold-expert & SM89 dense INT8 Tensor-Core GEMV optimizations: - Add 8-column Int2Group8PrepackedAvx2 in qmoe_shared_layout_int2_avx2.cc so all 8 columns of each 128-byte SM80 tile share one activation load while scaling weights before products. - Hoist activation loads outside step_col in Int3Group4Avx2 and implement a true 8-column Int3Group8Avx2 in qmoe_shared_layout_int3_avx2.cc. - Enable CanUseInt8TensorCore on SM89+ in fpA_intB_gemv/dispatcher.h and fix MSVC /WX C4459 variable shadowing in paged_attention_impl.cu. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…and skip zero-cold PCIe input copy
1. Add kOrtSessionOptionsMoeOffloadMinExpertWeight ("session.moe.offload.min_expert_weight", env fallback ORT_QMOE_MIN_EXPERT_WEIGHT):
- Prunes non-top-1 routed experts (rank >= 1) whose probability after SoftmaxTopK (and optional top-k normalization) is strictly below min_expert_weight, setting their routing scale to 0 and index to -2 while always preserving the top-1 expert (rank 0).
- Supports ORT_QMOE_SKIP_TAIL_COLD_ONLY=1 to retain GPU-resident (slot >= 0) tail experts during decode at zero wall-time cost while pruning only CPU-cold (slot < 0) tail experts, and ORT_QMOE_RENORM_AFTER_PRUNE=1 to optionally renormalize surviving weights.
- Skip zero-weight routes in QMoEOffloadMergeFinalizeCombineKernel and finalizeMoeRoutingKernel without reading uninitialized contribution rows.
2. In QMoEOffloadClassifyKernel, skip copying input_dev over PCIe into host_input_out when has_doorbell && cold_count == 0.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
a81fd9c to
29056d2
Compare
…default
Add kOrtSessionOptionsMoeOffloadSkipTailColdOnly ("session.moe.offload.skip_tail_cold_only",
default "1", overriding ORT_QMOE_SKIP_TAIL_COLD_ONLY) so that when
session.moe.offload.min_expert_weight > 0, decode retains GPU-resident (slot >= 0) tail
experts at zero wall-time cost and prunes only CPU-cold (slot < 0) tail experts by default.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- Keep QMoE warmup controlled solely by per-run controls in IsInQMoEWarmup - Preserve successful post-reclaim decode pinned reservation in EnsurePersistentDecodeState - Keep explicit device_expert_bytes as an upper bound and guard platform_free_vram with _WIN32 - Extend Int8TensorCoreGemvTest to SM 8.9+, report actual device SM, and assert tensor_core=1 - Fix clang-format comment alignment in moe_offload_plan.h Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Use CudaStream ownership to initialize valid library handles and release the async same-buffer copy fixtures safely. Pin SentencePiece 0.2.2 consistently for CI and transformer tools so its generated protobuf code remains compatible with protobuf 6.33.6. Validation: 3 CUDA self-copy tests, 7 offload policy tests, and 6 T5 generation tests passed; focused lint is clean. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Baiju Meswani (baijumeswani)
left a comment
There was a problem hiding this comment.
This PR improves QMoE CPU/GPU offload for Windows, updates memory planning and staging, widens the AVX2 INT2/INT3 CPU paths, enables the INT8 tensor-core path on SM89, and adds optional pruning of low-weight expert routes. These are useful directions for running larger models on limited GPU memory.
What looks good: the packed layouts and numerical behavior held up in extensive fresh CPU tests. The earlier warmup-control, duplicate pinned-reservation and SM89-test concerns are addressed. The ordinary explicit GPU-budget case is also fixed, although one boundary case remains below.
I do not think this is ready for approval yet. The six inline comments cover two memory-limit violations, two pruning configuration/path inconsistencies, and two repeatable CPU projection slowdowns. Please preserve explicit memory limits, honor an explicit zero pruning threshold, make pruning consistent across decode paths, and retain or improve the CPU paths that are faster in the measured configuration. Add regression tests for the limit and pruning cases, and repeat the CPU comparisons with production build settings.
Validation: two independent reviewers covered CPU math/performance and memory/planning, alongside CUDA/integration review. Fresh checks passed 16 CPU tests (one additional AVX512-only test skipped), 39 planning/store/retry tests, 56 production CUDA scenarios, and eight adapted INT8 tests. Targeted CUDA memory checks were clean, and native/plugin offload frontends compiled. The two Windows limit tests fail at this head and pass at the base and with small corrective changes in the test artifacts.
Limits: Windows planning was tested through exact-source simulations, not native WDDM execution. INT8 tests exercised pre-SM90 code on A100 with a review-only SM89 reporting adaptation, not actual SM89 hardware. CPU performance measurements are Linux/EPYC/GCC13 projection-helper timings, not full-model or Windows throughput. No full-model speed or quality result is claimed.
…090 QMoE offload optimizations - Preserve explicit host_expert_bytes cap and bounded file cache in QMoESessionController::PlanAndAdmitSession, shrinking within the cap when host RAM is smaller - Check mandatory fixed overhead against explicit device_expert_bytes cap before WDDM live budget clamping - Honor explicit session.moe.offload.min_expert_weight=0 over ORT_QMOE_MIN_EXPERT_WEIGHT in v4 planned mode - Apply cold-only tail pruning consistently across doorbell, callback, and all-cold decode paths before D2H copy and CPU scheduling - Retain short-lived activation loads in Int3Group4Avx2 and 4-column Int2Group4Avx2 path to avoid register spills, and integrate UnpackDequant8Int3WeightsAvx2 / UnpackDequant16Int3Avx512 / Int4LookupAvx512 - Integrate 16-expert integer-bank streaming wave default, >=2000-row stream_no_cpu_tail threshold, non-blocking promo_ready_ query, and MoeGemvTcInt4Kernel decode optimizations - Add C++ and Python regression tests for mandatory overhead cap rejection, low-RAM host cap enforcement, env override opt-out, and callback/doorbell cold-only pruning equivalence Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Baiju Meswani (baijumeswani)
left a comment
There was a problem hiding this comment.
Thanks for the updates. This PR improves QMoE CPU/GPU offload and memory planning, and adds optional pruning of low-weight expert routes. Five of the six earlier findings are now resolved: the explicit GPU limit, explicit zero pruning setting, callback/all-cold pruning behavior, and both CPU performance regressions.
The fixes look good in the focused checks. INT2 performance is back within measurement noise, and the tested INT3 projection is about 15% faster than the original base. Fresh validation passed 20 CPU tests (two additional AVX512-only tests skipped), 34 native policy/admission tests, 736 INT4 parity cases, 120 pruning-pipeline cases and 56 routing/finalization cases. Targeted CUDA memory checks were clean, and native/plugin offload frontends compiled.
I would still hold approval for the remaining host-memory issue in the inline comment. The planner now preserves the explicit host limit, but the backend constructor can disable the bounded cache afterward. Please preserve that admitted limit and add a regression covering the constructor and cache initialization, not only the planner.
The source-level Windows regression reproduces this issue and passes when the constructor override is removed. Windows behavior was simulated; native Windows, SM89 and AVX512 execution were not locally validated. No full-model speed or quality result is claimed.
…ion and working-set retention (#33935) ### Description Follow-up to #33916. Profiling QMoE CPU/GPU offload decode on an RTX 4090 (24 GB VRAM, ~30 GiB host RAM, Windows) identified several decode bottlenecks. This PR addresses them: 1. **AVX2 register spills (MSVC)** in `Int2Group4Avx2`, `Int3Group8Avx2`, `Int4Group8Avx2`, `Int4Group4Avx2`: captured activation arrays were spilled to the stack, and `RT = 3, 4` instantiations needed more than 16 YMM registers. Activations are now loaded with short-lived `_mm256_loadu_ps` folded into FMA, low/high halves are accumulated sequentially, and `RT <= 2`. Also adds fast finite BF16 scale rounding paths to `Int3LookupAvx2` and `RoundInt4WeightsAvx2`. 2. **Working-set eviction of hot cold experts**: dummy-token warmup faulted ~11.4 GiB of random experts into the working set, and `PrefetchVirtualMemory` on later turns evicted hot decode experts on low-RAM hosts. The warmup working set is now trimmed before the first live turn (Windows, < 48 GiB RAM), and `PrefetchVirtualMemory` is gated on available physical RAM. 3. **Pinned bounce starvation**: `SessionPromotionBudget` bounce slots are now sized to fit the remaining session pinned budget, so both bounce epochs stay allocated during decode. 4. **Eager `CpuWork` allocation / unneeded spin-waits**: scratch buffers are allocated lazily in `RunColdExperts`, and direct 2-pass `SpinPool` execution is used when `FileExpertStore` demand loading is inactive. 5. **2K context tier WDDM reserve**: reduce the extra WDDM reserve on < 28 GiB GPUs for the 2K tier (5,742 -> 6,573 resident GPU slots, decode hit rate 51.46% -> 58.71%). ### Results (RTX 4090, Qwen3.8-Flash-Next RC5, 3 reps, 64 new tokens) | Context | Metric | Before | After | | :--- | :--- | :---: | :---: | | Isolated 6-layer cold expert (H=2560, I=640) | Latency | 13.64 ms | 10.75 ms (-21%; INT4 -47%, INT3 -48%, INT2 -44% per expert) | | 2K | Resident GPU slots | 5,742 | 6,573 | | 2K | Decode GPU cache hit rate | 51.46% | 58.71% | | 2K | CPU prep time / step | 6.112 ms | 3.241 ms | | 2K | Median decode throughput | 11.86 tok/s | 23.07 tok/s | | 8K | CPU prep time / step | 8.337 ms | 5.410 ms | | 8K | Median decode throughput | 17.23 tok/s | 24.82 tok/s | | 8K | Peak VRAM | 23,659 MiB | 22,679 MiB | 8K outputs remain bit-exact (identical hashes) versus the baseline. ### Motivation and Context Improves decode throughput of QMoE CPU/GPU hybrid offload on consumer GPUs with limited VRAM and host RAM. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Summary
This PR fixes Windows WDDM and bounded-host-cache issues in QMoE CPU/GPU offload, adds tail-expert probability threshold pruning (
session.moe.offload.min_expert_weight), and optimizes prefill and decode throughput for mixed-precision (INT4/INT3/INT2) QMoE models on consumer GPUs (e.g., RTX 409024 GBVRAM +30 GiBRAM):FileExpertStore::InitializeMapped&QMoEOffload::InitializeBoundedHostCache(qmoe_offload.h,qmoe_offload.cc,qmoe_shared_layout.cc):external_file_pathandexternal_file_offsetfromMoeExpertSourceDescriptorintoQMoEOffload::InitializeBoundedHostCache().#elif defined(_WIN32)inFileExpertStore::InitializeMapped), accept committedMEM_MAPPED(orread_opt_enabled) mappings viaVirtualQuerywhensource.file_path == nullptr.qmoe_offload.cc):SessionPromotionBudget::HeldBytes()and computereserved = bounce_held + std::max(decode_held, reserve_target)inQMoEOffload::ComputeSharedso promotion bounce allocations (16 MiB) are not conflated with persistent decode headroom (24 MiB).QMoEOffload::EnsurePersistentDecodeState, release idle promotion bounce and landing-ring pinned staging whentry_reserve_decode_pinned()needs headroom.qmoe_session_controller.h):ullTotalPhys < host_expert_bytes + 24 GiB, automatically prefer direct OS file-backed mapping (use_bounded_file_cache = falseunlessORT_QMOE_FORCE_BOUNDED_HOST_CACHE=1), avoiding multi-GiB private heap commits and enabling 2-bank parallel pinned wave staging.num_layers >= 24,total_vram >= 16 GiB,total_full_gpu_bytes >= 8 GiB), automatically raise implicitdevice_working_bytesto560 MiBand implicitpinned_bytesto160 MiB, and adaptadmitted_gpu_expert_budgetagainstplatform_free_vram(DXGI local budget) andSessionBudgetAccountwithout spilling into WDDM non-local shared memory.EvictUploadedHostPages(VirtualUnlock) on WindowsMEM_MAPPEDviews.PrefetchVirtualMemory(qmoe_offload.cc):stream_rowsto8192acrossINT2,INT3, and full-budgetINT4layers.wave_experts(32 -> 16) inComputeSharedwhen needed soINT3layers fitstage_weight_banks = 2(full double-buffered CPU staging + PCIe H2D DMA + GPU Tensor-Core GEMM overlap).ORT_QMOE_WIN_PREFETCH_MMAP(default1on Windows) to all file-mapped integer-bank layers.session.moe.offload.min_expert_weight,qmoe_kernels.cu,moe_kernels.cu,qmoe_offload.cc):kOrtSessionOptionsMoeOffloadMinExpertWeight(session.moe.offload.min_expert_weight, env fallbackORT_QMOE_MIN_EXPERT_WEIGHT) to skip non-top-1 routed experts (rank >= 1) whose probability afterSoftmaxTopKis strictly belowmin_expert_weight(topk_scales[route] = 0.0f,topk_indices[route] = -2), always keeping at least the top-1 expert (rank == 0).ORT_QMOE_SKIP_TAIL_COLD_ONLY=1to retain GPU-resident (slot >= 0) tail experts during decode at zero wall-time cost while pruning only CPU-cold (slot < 0) tail experts insideQMoEOffloadClassifyKernel.QMoEOffloadMergeFinalizeCombineKernelandfinalizeMoeRoutingKernelwithout reading uninitialized contribution rows, and skipinput_dev -> host_input_outPCIe copies inQMoEOffloadClassifyKernelwhenhas_doorbell && cold_count == 0.INT2/INT3cold-expert &SM89denseINT8Tensor-Core GEMV optimizations (qmoe_shared_layout_int2_avx2.cc,qmoe_shared_layout_int3_avx2.cc,dispatcher.h,paged_attention_impl.cu):Int2Group8PrepackedAvx2so all 8 columns of each 128-byte SM80 interleaved tile share one activation load while scaling weights before products.step_colinInt3Group4Avx2and implement a true 8-columnInt3Group8Avx2.CanUseInt8TensorCoreonparams.device_sm >= 89(SM89) and fix MSVC/WXC4459variable shadowing inpaged_attention_impl.cu.Testing
test_qmoe_offload_v4.pytargeted unit tests: PASS.qwen_flash_next_int4_int8_rc5_contrib_dflash2_qmoe_2b375_x02_y14_z32_dp23_b64) benchmarks on RTX 4090 (24 GBVRAM) + Intel i9-13900 (30 GiBRAM):2K(2,048tokens): TTFT8.92 s(7.55 satmin_expert_weight=0.08), decode12.15–23.43 tok/s(24.63–29.63 tok/swithmin_expert_weight=0.065–0.08).8K(8,192tokens): TTFT14.42–15.32 s(10.84–12.30 swithmin_expert_weight=0.065–0.08), decode23.04 tok/smedian /29.49 tok/speak.32K(32,768tokens): TTFT71.71–73.16 s, decode8.22–9.48 tok/s.128K(131,072tokens): TTFT536.21–579.64 s, decode7.42–15.27 tok/s, peak VRAM21,345 MiB.shard 0 of 8,max-new 8192):min_expert_weight=0): 86/100 (86.0%),0truncated,27.39 tok/saggregate decode (3,196 swall time).min_expert_weight=0.065,skip_cold_only=1): 84/100 (84.0%),0truncated,34.70 tok/saggregate decode (43.21 tok/smax,2,099 swall time;43.10 tok/savg /53.42 tok/smax at0.080on 20-item subset).