Repository navigation
Support direct external initializer buffers in AddExternalInitializers - #33877
Conversation
…rnalInitializers Honor session.use_external_initializer_file_buffers_directly in Graph::InjectExternalInitializedTensors so naturally aligned, little-endian caller buffers (such as mmap'd GGUF weight tensors) are borrowed in place without a duplicate CPU allocation and copy. Also allow external initializers that already reside on the planned device to be used in place during SaveInitializedTensors without duplicate arena planning or device-to-device copies. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
There was a problem hiding this comment.
🟡 Changes recommended
Non-owning device buffers are borrowed without opt-in, risking dangling pointers, while optimizer reads use host memcpy on device memory and can crash session initialization.
4 open issues
What changed in this PR
This PR extends external-initializer loading to reduce duplicate weight allocations during session initialization, particularly for large memory-mapped models.
Changes:
- Passes the direct-buffer option into graph injection to borrow eligible caller-owned CPU buffers.
- Adds device-buffer handling and skips redundant allocation planning when initializers already reside on the planned device.
- Updates configuration documentation and adds framework and public API tests.
| File | Description |
|---|---|
onnxruntime/test/shared_lib/test_model_loading.cc |
Adds a public API test with direct-buffer mode enabled. |
onnxruntime/test/framework/inference_session_test.cc |
Checks default copying, opt-in borrowing, and inference results. |
onnxruntime/core/session/inference_session.cc |
Passes the direct-buffer configuration into graph injection. |
onnxruntime/core/graph/graph.cc |
Adds borrowing paths for eligible CPU and device buffers. |
onnxruntime/core/framework/tensorprotoutils.cc |
Uses in-memory references for small nonempty device tensors. |
onnxruntime/core/framework/session_state_utils.cc |
Skips allocation planning and reuses initializers on the planned device. |
include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h |
Expands the configuration comment to cover AddExternalInitializers. |
include/onnxruntime/core/graph/graph.h |
Adds the optional direct-buffer parameter. |
docs/Model_Loading_Performance.md |
Documents buffer borrowing and device-initializer reuse. |
Review effort: Balanced
Justin Chu (justinchuby)
left a comment
There was a problem hiding this comment.
Can this be set to true by default?
There was a problem hiding this comment.
🟡 Changes recommended
Critical device-buffer handling can copy CUDA pointers as host memory during initialization and borrow caller-owned buffers without opt-in, risking crashes and use-after-free.
8 open issues
Document borrowed-buffer lifetime requirements across APIs · New Stage device-backed initializers for optimization and shape inference · New ConvAdd fusion reads device initializers with host memcpy Device buffers are borrowed even when direct use is disabled Add CUDA tests for device reuse and cross-device transfers · New Test the misaligned-buffer fallback in direct mode · New Test initializer is too small to exercise buffer borrowing External float buffers are misread on big-endian hosts
Review effort: Max
Preserve the default caller-buffer lifetime contract and keep graph initializers host-readable during optimization and shape inference. Document opt-in CPU borrowing and cover device, endian, and misaligned buffers with focused regressions. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
|
Justin Chu (@justinchuby) I kept borrowing opt-in to preserve the existing caller-buffer lifetime contract. Applications can currently release non-owning initializer buffers after session creation; enabling borrowing by default would silently require those buffers to remain unchanged and alive for the lifetime of every consuming session. CPU OrtValues that own their buffers already retain ownership without requiring this option, so safe owned-buffer sharing remains available. The update in dd1ccfa also always stages device initializers into owned CPU storage before graph processing instead of introducing device-buffer borrowing. |
There was a problem hiding this comment.
🟡 Changes recommended
Critical memory-safety issues remain: borrowed weights can trigger unaligned-access undefined behavior, and staging non-contiguous device initializers can overrun host buffers.
3 open issues
8 resolved since last review
Stage device-backed initializers for optimization and shape inference Document borrowed-buffer lifetime requirements across APIs ConvAdd fusion reads device initializers with host memcpy Device buffers are borrowed even when direct use is disabled Test the misaligned-buffer fallback in direct mode Add CUDA tests for device reuse and cross-device transfers Test initializer is too small to exercise buffer borrowing External float buffers are misread on big-endian hosts
Review effort: Max
Track baseline workers by Linux thread ID instead of an absolute count so threads lingering after join cannot inflate expected worker counts. Add deterministic coverage for a baseline worker that exits. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Reject strided external initializers before CPU injection or device-to-host staging, and use an unaligned-safe quantized load so borrowed byte buffers remain valid for MatMulNBits with g_idx. Correct public API ownership documentation and add CPU/device stride and direct tensor/file buffer regressions. Validated 19 focused tests in a CPU training build with alignment UBSan, plus 10 repetitions of each new regression. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…block-packed weights (#33953) ## Description Adds three `com.microsoft` contrib operators—`GatherPacked` (CPU and CUDA), `MatMulPacked` (CUDA), and `MoEPacked` (CUDA)—that consume 2D and 3D block-packed quantized weight blobs (such as GGUF `IQ2_XXS`, `IQ2_XS`, `IQ2_S`, `IQ3_XXS`, `IQ4_NL`, `IQ4_XS`, `Q2_0`, `Q4_0`, `Q5_0`, `Q5_K`, `Q6_K`, `Q8_0`, `F16`, `BF16`, `F32`) directly in their native block layout without offline weight repacking or separate scale tensors. Direct external initializer buffers (#33877) avoid redundant host-side copies while loading `.gguf` shards. CUDA execution still copies CPU-backed weights to device memory and requires device storage for the packed weights; it does not execute zero-copy from host-backed shards. ## Summary of Changes ### Contrib Operator Schemas & Registration | File | Change | |------|--------| | `onnxruntime/core/graph/contrib_ops/ms_opset.h` | Declare `GatherPacked`, `MatMulPacked`, and `MoEPacked` in `com.microsoft` opset | | `onnxruntime/core/graph/contrib_ops/contrib_defs.cc` | Register schemas and shape inference for `GatherPacked`, `MatMulPacked` (`A, B -> Y` with `K, N, packed_type`), and `MoEPacked` | | `onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc` | Register CPU `GatherPacked` kernel | | `onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc` | Register CUDA `GatherPacked`, `MatMulPacked`, and `MoEPacked` kernels | ### Block-Packed Quantization & Kernel Implementations | File | Change | |------|--------| | `onnxruntime/contrib_ops/cpu/quantization/gguf_quant_tables.h` | Shared block-quantization lookup tables (`iq2xxs_grid`, `iq2xs_grid`, `iq2s_grid`, `iq3xxs_grid`, `iq3s_grid`), block geometry helper (`GgmlRowBytes`), and CPU row dequantizer | | `onnxruntime/contrib_ops/cpu/quantization/gather_packed.cc` | Multi-threaded CPU `GatherPacked` implementation | | `onnxruntime/contrib_ops/cuda/quantization/gguf_quant_cuda.cuh` | CUDA sub-block decoders (`DecodeGgufSubBlock8Impl`, `DecodeGgufSubBlock8SmemLutTpl`, `DecodeGgufSubBlock16Impl`), register-only `Iq2xsSignsDev` / `Iq4NlValueDev` helpers, and Tensor Core `MmaBf16` wrapper | | `onnxruntime/contrib_ops/cuda/quantization/gather_packed.{h,cc,cu}` | CUDA `GatherPacked` kernel | | `onnxruntime/contrib_ops/cuda/quantization/matmul_packed.{h,cc,cu}` | CUDA `MatMulPacked` kernel with warp-K-split Tensor Core (`mma.sync.aligned.m16n8k16`) and warp GEMV paths for decode (`M <= 8`), bounded tiled dequantization + cuBLAS GEMM for prefill (`M > 8`), and FP32 decoded weights for FP16 prefill when the packed format is not F16 | | `onnxruntime/contrib_ops/cuda/moe/moe_packed.{h,cc,cu}` | CUDA `MoEPacked` kernel with fused top-k softmax routing, parallel warp-scan tile bucketing (`8` routes/tile in decode, `32` routes/tile in prefill), warp-K-split Tensor Core FC1+SwiGLU in decode, 4-subtile Tensor Core FC1/FC2 in prefill, and 128-bit vectorized finalize | | `onnxruntime/test/python/transformers/test_gguf_packed_ops.py` | Unit tests verifying `GatherPacked` (CPU vs CUDA bit-exactness across 15 formats), `MatMulPacked` (`M = 1, 7, 8, 32`), and `MoEPacked` (`M = 4, 48`) | ## Testing - Additional regression validation on NVIDIA A100 with CUDA 13: **30 tests and 757 subtests pass** both with the normal SM80/SM90a build and with the packed-op kernels compiled as compute_75-only PTX, exercising BF16 SIMT fallback. - Configured-device capability probing and restoration were verified on both success and failure; focused `lintrunner` checks pass. - Unit tests: ```bash python -m unittest onnxruntime/test/python/transformers/test_gguf_packed_ops.py -v ``` - End-to-end validation on `Qwen3.8-Flash-Next-GGUF` (`UD-IQ1_S`, `UD-IQ2_XS`, `UD-IQ3_XXS`) on NVIDIA H200: - `engine_behaviours.py`: **27 / 27 PASS** across all three quantization variants. - Prefill (TTFT): **0.048–0.052 s** at 128 tokens, **0.222–0.244 s** at 2K tokens (**8,400–9,226 tok/s**), and **0.851–0.944 s** at 8K tokens (**8,682–9,622 tok/s**, **17.1x–18.5x** speedup over baseline SIMT). - Speculative decode (`DFlash2 S=8`): **20.0–20.9 ms/cycle** (**183–269 tok/s** median, up to **359 tok/s** peak). - MMLU-Pro (200 items, `max_new=8192`): **87.50%** (`UD-IQ1_S`), **85.50%** (`UD-IQ2_XS`), **86.00%** (`UD-IQ3_XXS`). ## Checklist - [x] Tests added/updated - [x] No breaking changes (or documented in description) - [x] Documentation updated (if applicable) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>



Description
Extend
session.use_external_initializer_file_buffers_directlytoAddExternalInitializersso applications can opt into borrowing eligible non-owning CPU weight buffers instead of allocating and copying redundant host storage. Borrowing requires naturally aligned buffers larger than 127 bytes on little-endian hosts; default copying behavior remains unchanged.Device initializers are always copied into owned CPU staging memory before graph injection, optimization, or shape inference, using the registered execution provider's data-transfer implementation. This PR does not enable device-buffer borrowing or same-device GPU zero-copy reuse.
Motivation and Context
Large memory-mapped models otherwise incur a redundant host allocation and copy when their non-owning weight views are supplied through
AddExternalInitializers. Keeping borrowing opt-in preserves the existing guarantee that copied caller buffers can be released after session creation. Borrowed CPU buffers must instead remain unchanged and alive until all sessions using them are released.Summary of Changes
AddExternalInitializers, while preserving existing ownership retention for CPU buffers owned by theirOrtValue.Add, Level1Conv/Addfusion, small and large device buffers, source-buffer overwrite/release, and CPU-planned shape initializers.Testing
Validated on Linux/H200 with the existing CUDA 13.0 / cuDNN 9.23 build configuration.
cmake --build build/cu130/Release --target onnxruntime_test_all onnxruntime_shared_lib_test --parallel 12: passed.build/cu130/Release:./onnxruntime_test_all --gtest_filter='GraphTest.InjectExternalInitializedTensors*:InferenceSessionTests.*ExternalInitializer*:InferenceSessionTests.ExternalDeviceInitializers*': 11 passed, none skipped.build/cu130/Release:./onnxruntime_shared_lib_test --gtest_filter='*ExternalInitializer*': 8 passed, none skipped.lintrunner --revision HEADbefore committing the feedback fixes: passed.git diff --check: passed.The device-buffer aliasing regression failed against the original PR and passes after the fix. Big-endian execution was not validated locally; the original GGUF end-to-end workload was not rerun.
Checklist