Skip to content

Support direct external initializer buffers in AddExternalInitializers - #33877

Merged
Tianlei Wu (tianleiwu) merged 5 commits into
mainfrom
tlwu/20261104/gguf-bolt
Nov 8, 2026
Merged

Tianlei Wu (tianleiwu) merged 5 commits into
mainfrom
tlwu/20261104/gguf-bolt

Conversation

@tianleiwu

@tianleiwu Tianlei Wu (tianleiwu) commented Nov 4, 2026 •

Copy link
Copy Markdown
Contributor

Description

Extend session.use_external_initializer_file_buffers_directly to AddExternalInitializers so 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

  • Honor the direct-buffer option for eligible CPU tensors supplied through AddExternalInitializers, while preserving existing ownership retention for CPU buffers owned by their OrtValue.
  • Stage device tensors on CPU before graph processing, reject graph injection of device pointers, and report missing data-transfer support explicitly.
  • Remove the proposed small-device address-tag exception and device-specific allocation-planning changes.
  • Document copying versus borrowing and buffer lifetime requirements in the C API, configuration keys, and model-loading guide.
  • Add CUDA regressions for scalar Add, Level1 Conv/Add fusion, small and large device buffers, source-buffer overwrite/release, and CPU-planned shape initializers.
  • Correct the framework test's native-endian backing storage, add a graph-level misaligned-buffer regression, and verify public API buffer reuse with a 256-byte initializer observed inside a test kernel.

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.
  • From build/cu130/Release: ./onnxruntime_test_all --gtest_filter='GraphTest.InjectExternalInitializedTensors*:InferenceSessionTests.*ExternalInitializer*:InferenceSessionTests.ExternalDeviceInitializers*': 11 passed, none skipped.
  • From build/cu130/Release: ./onnxruntime_shared_lib_test --gtest_filter='*ExternalInitializer*': 8 passed, none skipped.
  • lintrunner --revision HEAD before 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

  • Tests added/updated
  • Default CPU copying and caller-buffer lifetime behavior preserved
  • Documentation updated

…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>

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

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

Comment thread onnxruntime/core/graph/graph.cc Resolved
Comment thread onnxruntime/core/graph/graph.cc Resolved
@tianleiwu
Tianlei Wu (tianleiwu) marked this pull request as draft November 4, 2026 17:33

@justinchuby Justin Chu (justinchuby) 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.

Can this be set to true by default?

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.

Comment thread onnxruntime/core/graph/graph.cc Resolved
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>
@tianleiwu

Copy link
Copy Markdown
Contributor Author

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.

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.

Comment thread onnxruntime/core/graph/graph.cc Resolved
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>

@rui-ren Rui Ren (rui-ren) 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.

LGTM, thanks!

@tianleiwu
Tianlei Wu (tianleiwu) merged commit 3875bc0 into main Nov 8, 2026
102 checks passed
@tianleiwu
Tianlei Wu (tianleiwu) deleted the tlwu/20261104/gguf-bolt branch November 8, 2026 20:14
Tianlei Wu (tianleiwu) added a commit that referenced this pull request Nov 9, 2026
…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>
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.

4 participants