Skip to content

ggml: use dynamic allocation for split graph inputs - #22789

Merged
ggerganov merged 5 commits into
ggml-org:masterfrom
AgoraPete:fix/dynamic-split-inputs
Aug 3, 2026
Merged

ggml: use dynamic allocation for split graph inputs#22789
ggerganov merged 5 commits into
ggml-org:masterfrom
AgoraPete:fix/dynamic-split-inputs

Conversation

@AgoraPete

Copy link
Copy Markdown
Contributor

Overview

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically allocated buffers in the backend scheduler. This fixes crashes when loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on multi-backend setups where graph splits exceed 30 input tensors. Done with the lightest touch possible to the fewest files to fix the issue.

Additional information

  • split->inputs: dynamic array with grow-on-demand
  • sched->graph_inputs: dynamic array with grow-on-demand
  • graph_size calculation now uses actual input count instead of fixed constant

Requirements

  • I have read and agree with the contributing guidelines
  • AI usage disclosure: Yes - requirement was surfaced during local debugging of Gemma4 Draft MTP on edge device. The model was Big Pickle if you're interested.

AgoraPete added 2 commits May 6, 2026 20:30
Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant
@github-actions github-actions Bot added the ggml changes relating to the ggml tensor library for machine learning label May 7, 2026
@ggerganov ggerganov self-assigned this May 11, 2026
@opticblu

Copy link
Copy Markdown

Possible HIP/ROCm issue with --no-mmap / --no-host after dynamic split inputs patch

First, thank you for working on this patch.

I’m seeing two related failure modes on ROCm/HIP when changing the memory options (using fix/dynamic-split-inputs), so I wanted to report them

Environment:

  • OS: Debian 13 / Trixie
  • Kernel: 6.19.11+deb13-amd64
  • GPUs: 2x AMD Radeon AI PRO R9700, gfx1201
  • ROCm: 7.2.3
  • llama.cpp build: HIP/ROCm
  • Build flags: GGML_HIP=ON, GPU_TARGETS=gfx1201, GGML_HIP_ROCWMMA_FATTN=ON, CMAKE_BUILD_TYPE=Release
  • Tested against a build containing 72e90b92c ggml: use dynamic allocation for split graph inputs

The default mmap path appears to work for me.

When I add --no-mmap, generation starts normally but later crashes around token ~1000 with an illegal memory access in the HIP backend.

Example command:

ROCR_VISIBLE_DEVICES=0,1 /opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/llama-server -hf unsloth/Qwen3.6-35B-A3B-GGUF:Q8_K_XL --host 0.0.0.0 --port 8080 --gpu-layers all --jinja --api-key REDACT -fa on -ts 1,1 -kvu -kvo --no-mmap
/opt/src/llama.cpp-dynamic-split-inputs/ggml/src/ggml-cuda/ggml-cuda.cu:102: ROCm error
ROCm error: an illegal memory access was encountered
  current device: -1, in function ggml_backend_cuda_event_wait at /opt/src/llama.cpp-dynamic-split-inputs/ggml/src/ggml-cuda/ggml-cuda.cu:4523
  hipStreamWaitEvent(cuda_ctx->stream(), (hipEvent_t)event->context, 0)
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-base.so.0(+0x172a5) [0x7f26c5dd72a5]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-base.so.0(ggml_print_backtrace+0x1df) [0x7f26c5dd767f]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-base.so.0(ggml_abort+0x11e) [0x7f26c5dd780e]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-hip.so.0(+0x359ecf2) [0x7f26c457ecf2]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-hip.so.0(+0x35a5b20) [0x7f26c4585b20]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libggml-base.so.0(ggml_backend_sched_graph_compute_async+0x333) [0x7f26c5df3b13]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libllama.so.0(_ZN13llama_context13graph_computeEP11ggml_cgraphb+0xa1) [0x7f26c5f3f491]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libllama.so.0(_ZN13llama_context14process_ubatchERK12llama_ubatch14llm_graph_typeP22llama_memory_context_iR11ggml_status+0xe4) [0x7f26c5f41fd4]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libllama.so.0(_ZN13llama_context6decodeERK11llama_batch+0x34f) [0x7f26c5f47c8f]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/libllama.so.0(llama_decode+0xb) [0x7f26c5f496bb]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/llama-server(+0x10b1a2) [0x5558ff1711a2]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/llama-server(+0x19e8b1) [0x5558ff2048b1]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/llama-server(+0x691f5) [0x5558ff0cf1f5]
/lib/x86_64-linux-gnu/libc.so.6(+0x29ca8) [0x7f26c585eca8]
/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x85) [0x7f26c585ed65]
/opt/src/llama.cpp-dynamic-split-inputs/build-r9700-hip/bin/llama-server(+0x698f1) [0x5558ff0cf8f1]
Aborted

I also found a possibly related failure mode with --no-host.

Even without --no-mmap, adding --no-host does not crash immediately, but generation becomes corrupted and the model starts repeating itself badly after some output.

Current behavior matrix:

default mmap path: appears to work
--no-mmap: ROCm illegal memory access around token ~1000
--no-host: generation corruption / severe repetition after some output
--no-mmap --no-host: not yet separately isolated
I do not know yet whether this is specific to ROCm/HIP, tensor split across two GPUs, or the interaction with the new dynamic split input allocation path. But since the default path works and these memory options change the behavior substantially, I wanted to report it while the patch is still under review.

Can anyone else confirm this? Thanks

@nbCloud91

Copy link
Copy Markdown

Confirming both the bug in #24132 and this PR's fix on a low-VRAM CUDA setup — the hardware class mentioned above as having no current options.

Setup: Quadro T1000 4GB (Turing), CUDA backend, Linux (NixOS); b10063 + this patch (applies cleanly); gemma-4-E2B-it-Q4_0 + mmproj from the ggml-org GGUF repo, via llama-mtmd-cli --jinja.

Without the patch: -ngl 99 OOMs (expected on 4GB — the 540 MiB compute buffer doesn't fit next to full weights); any partial offload, including the default -fit path, crashes with GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS) failed at ggml-backend.cpp:1367.

With the patch: -fit completes ("successfully fit params to free device memory", 0.82 s), offloads 36/36 layers with weights CPU-mapped, peak VRAM 2888 MiB of 4096 MiB, and runs stably:

text gen (48-tok run, llama_perf) prompt eval
patched, -fit default 58.56 t/s 165.65 t/s
same build, -ngl 0 (CPU) 17.93 t/s 228.57 t/s

Vision (image describe) and audio (8 s speech transcription) outputs are consistent with the CPU-only baseline — no quality change observed.

Re the ROCm --no-mmap/--no-host report above: CUDA with default memory settings (mmap enabled) shows no issue here; I did not test --no-mmap/--no-host. One observation that may help isolate it: besides the dynamic arrays, the patch also changes graph_size from the GGML_SCHED_MAX_SPLIT_INPUTS worst-case to actual input counts — a smaller reservation than before, which might be worth separating from the dynamic-array change if the ROCm issue reproduces.

Testing was done with AI assistance (Claude).

@LeConsulat2

Copy link
Copy Markdown

I can reproduce this assertion consistently with the official Windows Vulkan build on two different Intel integrated GPUs.

The failure occurs only with partial CPU/GPU offload. CPU-only and full GPU offload both work correctly.

Common setup

  • llama.cpp build: b9585 (d73cd0767)
  • Binary: official Windows Vulkan release, llama-b9585-bin-win-vulkan-x64.zip
  • Compiler reported by the binary: Clang 20.1.8, Windows x86_64
  • ZIP SHA-256: af6b1b94377b9f78dbb2285b878fb696d36766391499d65e055ecd622b69018a
  • Backend: Vulkan
  • Model: gemma-4-E2B-it-qat-UD-Q4_K_XL.gguf from the unsloth GGUF repository
  • The model reports n_layer = 35; llama.cpp counts the repeating layers plus the output layer as 36 offloadable layers.

Command used:

.\llama-server.exe `
  -m .\gemma-4-E2B-it-qat-UD-Q4_K_XL.gguf `
  --host 127.0.0.1 `
  --port 8081 `
  -c 32768 `
  -t <N> `
  -tb <N> `
  -np 1 `
  --no-ui `
  --cache-prompt `
  --cache-reuse 256 `
  --reasoning auto `
  --reasoning-budget 3072 `
  -ngl <0|18|999> `
  -dev Vulkan0

The thread count was adjusted to match each machine's logical processor count. Everything else was identical.

Result on both machines

  • -ngl 0: works
  • -ngl 999: works
  • -ngl 18: consistently crashes before the server reaches listening

-ngl 18 offloads 18 of the 36 offloadable layers reported by llama.cpp, creating a genuine CPU/GPU split.

The exact assertion is:

D:/a/llama.cpp/llama.cpp/ggml/src/ggml-backend.cpp:1367:
GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS) failed

Machine A

  • CPU: Intel Core i5-1135G7, 4 cores / 8 threads
  • GPU: Intel Iris Xe Graphics
  • Operating system: Windows 11
  • Vulkan driver: 32.0.101.6556
  • Threads: -t 8 -tb 8
  • Reported Vulkan memory before the failure: 7,238 MiB free

Results:

  • -ngl 0: 44.15 tok/s prefill, 11.62 tok/s decode
  • -ngl 999: 16.54 tok/s prefill, 8.14 tok/s decode
  • -ngl 18: assertion failure

Crash-adjacent log:

0.01.228.743 D sched_reserve: reserving full memory module
0.01.228.769 D sched_reserve: worst-case: n_tokens = 512, n_seqs = 1, n_outputs = 1
0.01.228.770 D graph_reserve: reserving a graph for ubatch with n_tokens = 1, n_seqs = 1, n_outputs = 1
0.01.230.229 I sched_reserve: Flash Attention was auto, set to enabled
0.01.230.243 I sched_reserve: resolving fused Gated Delta Net support:
0.01.230.244 D graph_reserve: reserving a graph for ubatch with n_tokens = 1, n_seqs = 1, n_outputs = 1
0.01.231.329 I sched_reserve: fused Gated Delta Net (autoregressive) enabled
0.01.231.335 D graph_reserve: reserving a graph for ubatch with n_tokens = 16, n_seqs = 1, n_outputs = 16
0.01.233.019 I sched_reserve: fused Gated Delta Net (chunked) enabled
0.01.233.028 D graph_reserve: reserving a graph for ubatch with n_tokens = 512, n_seqs = 1, n_outputs = 1
D:/a/llama.cpp/llama.cpp/ggml/src/ggml-backend.cpp:1367:
GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS) failed

This does not present as an out-of-memory failure. The process aborts specifically on GGML_SCHED_MAX_SPLIT_INPUTS, with 7,238 MiB of Vulkan device memory reported free immediately beforehand.

Machine B

  • CPU: Intel Core 7 240H, 10 cores / 16 threads
  • Integrated GPU: Intel Graphics
  • Discrete GPU also present: NVIDIA GeForce RTX 4050 Laptop GPU
  • Operating system: Windows 11
  • Intel Vulkan driver: 32.0.101.7076
  • Threads: -t 16 -tb 16

Two Vulkan devices were detected:

Vulkan0 : Intel(R) Graphics
Vulkan1 : NVIDIA GeForce RTX 4050 Laptop GPU

All tests were explicitly pinned to Vulkan0. The RTX 4050 was not used.

Results:

  • -ngl 0: works, 0/36 layers offloaded
  • -ngl 999: works, 36/36 layers offloaded to Vulkan0
  • -ngl 18: assertion failure

Crash-adjacent log:

0.00.144.675 I device_info:
0.00.146.112 I   - Vulkan0 : Intel(R) Graphics (16181 MiB, 15413 MiB free)
0.00.147.344 I   - Vulkan1 : NVIDIA GeForce RTX 4050 Laptop GPU (5921 MiB, 5153 MiB free)
0.00.147.384 I system_info: n_threads = 16 (n_threads_batch = 16) / 16
0.00.405.785 I llama_prepare_model_devices: using device Vulkan0 (Intel(R) Graphics) - 15413 MiB free
0.00.660.294 I print_info: n_layer = 35
0.00.660.295 I print_info: n_layer_all = 35
0.00.674.965 D load_tensors: layer 18 assigned to device Vulkan0, is_swa = 1
0.00.674.965 D load_tensors: layer 19 assigned to device Vulkan0, is_swa = 0
...
0.00.674.971 D load_tensors: layer 34 assigned to device Vulkan0, is_swa = 0
0.00.674.971 D load_tensors: layer 35 assigned to device Vulkan0, is_swa = 0
D:/a/llama.cpp/llama.cpp/ggml/src/ggml-backend.cpp:1367:
GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS) failed

This also does not present as an out-of-memory failure. The same assertion occurs with 15,413 MiB of Vulkan device memory reported free.

Observations

Both machines hit the same assertion in the same source file and line, using the same model and the same partial-offload setting.

The machines use different Intel GPU generations and materially different driver versions, but the result is identical:

CPU only       (-ngl 0)   → works
Partial split  (-ngl 18)  → assertion failure
Full offload   (-ngl 999) → works

This makes a hardware-specific or single-driver-specific cause less likely and points toward the shared graph-splitting path used when the model spans both CPU and GPU backends.

The failure pattern appears consistent with PR #22789, which replaces the fixed GGML_SCHED_MAX_SPLIT_INPUTS storage with dynamically allocated split-input storage.

I have not yet tested a Vulkan build containing PR #22789, so I cannot confirm that the patch resolves this on Vulkan. This report only confirms that the same fixed-size split-input assertion is reproducible on Vulkan across two separate Intel integrated GPUs.

The current workaround is to avoid partial offload and use either -ngl 0 or -ngl 999.

Thanks,
Jonathan

@ggerganov
ggerganov merged commit dbadb68 into ggml-org:master Aug 3, 2026
24 of 26 checks passed
@AgoraPete

AgoraPete commented Aug 3, 2026 via email

Copy link
Copy Markdown
Contributor Author

@thomasjfox

Copy link
Copy Markdown
Contributor

Just wanted to report: The commit fixes loading DSv4 Flash + DSpark draft model on dual RTX 6000 Pro Max-Q.

No manual raising of GGML_SCHED_MAX_SPLIT_INPUTS needed anymore.

Thanks a lot!

smalinin pushed a commit to smalinin/llama.cpp that referenced this pull request Aug 4, 2026
* ggml: use dynamic allocation for split graph inputs

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant

* cont : clean-up

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Nathanw1014 added a commit to Nathanw1014/llama.cpp that referenced this pull request Aug 9, 2026
ggml-org#22789 replaced the fixed 30-entry split input array with a growable one and,
in the same edit, changed the split-cutting heuristic from the constant to
split->inputs_capacity:

    -   if (split->n_inputs == GGML_SCHED_MAX_SPLIT_INPUTS) {
    +   if (split->n_inputs >= split->inputs_capacity) {

inputs_capacity starts at GGML_SCHED_MAX_SPLIT_INPUTS but doubles on demand and
is never reset for the life of the sched, so once a split slot grows, the
scheduler stops cutting there and the cut point ratchets up for every later
graph build. Longer splits mean every cross-backend input copy is materialised
at the split's start and stays live to its last use inside it, which raises the
peak the compute-buffer allocator has to cover - n_copies times over under
pipeline parallelism. Only multi-backend configurations can reach this.

Keep the growable array, which is what fixes the original >30-input assert, and
cut on the constant again as before ggml-org#22789. >= rather than == so the check keeps
firing for splits that did have to grow.

DeepSeek-V4-Flash UD-IQ3_XXS, gfx1151, -c 400000 -ub 2048 -fa 1 --fit off:
Vulkan0 compute buffer 4714.00 MiB and 9157 graph nodes, byte-identical to the
unpatched tree, and neither run grows a split past 30 inputs. Expected - one
Vulkan device plus the CPU backend cannot exercise the path on this box. The
reported case is 3 devices with pipeline parallelism.

test-backend-ops -o FLASH_ATTN_EXT, run alone on gfx1151: 13257/13295 on both
this and the unpatched tree, with the same 38 failing cases (identical case
list, all type_K=q8_0 prec=def kv_view=1). Pre-existing on the branch, not
touched by this change.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
satindergrewal pushed a commit to satindergrewal/llama.cpp that referenced this pull request Aug 11, 2026
* ggml: use dynamic allocation for split graph inputs

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant

* cont : clean-up

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
satindergrewal pushed a commit to satindergrewal/llama.cpp that referenced this pull request Aug 12, 2026
* ggml: use dynamic allocation for split graph inputs

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant

* cont : clean-up

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
mndodd added a commit to mndodd/llama.cpp that referenced this pull request Aug 12, 2026
74 upstream commits, 14 in our paths. Two textual conflicts, both expected,
plus ONE silent-drift resolution that did NOT conflict -- which is the whole
reason upstream-check.sh warns about auto-merges.

☠☠ fattn.cpp -- THE ONE THAT AUTO-MERGED CLEAN (upstream 66fa168 / ggml-org#25874)
  ggml-org#25874 extends oneDNN SDPA to non-F16 KV (Q4_0..Q8_0) and inserts its gate
  ABOVE the MKL gate -- which our 0801 sync had already deferred to MMA for
  exactly this reason. Its envelope (Q->ne[1] >= 32, K->ne[1] >= 1024,
  quantized KV now accepted) matches our deploy prefill exactly: q8_0 KV,
  ubatch 2048, any real depth. Verbatim it would have silently swapped the
  measured MMA kernel (2.87x served prefill at 43k) for an unmeasured one and
  staged the whole q8_0 KV cache to dense F16 first.
  ⇒ added the SAME conjunct, same wording, as the MKL gate below it:
    !ggml_sycl_fattn_mma_supported(dst). oneDNN keeps its full envelope
    everywhere MMA declines. A/B reachable both ways without a rebuild
    (GGML_SYCL_FATTN_MMA=0 / GGML_SYCL_FA_ONEDNN=0 -- the latter is checked
    inside _supported(), so it still governs the new early return).
  This is a precedence choice, not a revert, and it is UNMEASURED on this box.
  B59 is the finding that prices it.

  The three conflicts git DID flag in this file were cosmetic (debug kname
  lines + upstream re-calling get_best_fattn_kernel where we hoist it). Kept
  ours, added upstream's ONEDNN kname lines.

concat.cpp -- upstream 6c8dcaa (ggml-org#25852) is a duplicate of our own fix
  Kept OURS: strict superset (same launch geometry + the GGML_SYCL_CONCAT_WG
  door + the o[dim] hoist + the i64 loop var). Ours deliberately omits their
  WARP_SIZE floor, which would make WG=1 unreachable and destroy the
  one-binary positive control for the 8.36x launch fix. Noted in the source
  so the next sync does not resolve toward upstream.

Assessed, no action needed:
  272700b (ggml-org#26105) iGPU classification -- NULL here, B70 is discrete and
    still reports TYPE_GPU. Relevant later for PVC/Max-1100 and B51.
  dbadb68 (ggml-org#22789) dynamic split-graph inputs -- mechanical, auto-merged
    beside our GGML_SCHED_HANDOFF_CENSUS.
  596a579 (ggml-org#25784) DeepseekV4 MTP + DSpark -- ⚠ changes SHARED MTP
    plumbing (n_embd_out vs n_embd_inp in llama_context::decode,
    set_embeddings_layer_inp bounds, embeddings_layer_inp sized n_layer+1).
    We run MTP every step. NOT yet gated -- see the build/gate that follows.

☠ NOT YET BUILT, NOT YET GATED. Every absolute in f306 predates this merge.
fewtarius added a commit to fewtarius/CachyLLama that referenced this pull request Aug 13, 2026
…y (nathanw1014)

ggml-org#22789 replaced the fixed 30-entry split input array with a growable one and
in the same edit changed the split-cutting heuristic from the constant to
split->inputs_capacity. inputs_capacity starts at GGML_SCHED_MAX_SPLIT_INPUTS
but doubles on demand and is never reset, so once a split slot grows, the
scheduler stops cutting there and the cut point ratchets up for every later
graph build. Longer splits mean every cross-backend input copy is materialised
at the split's start and stays live to its last use inside it, which raises the
peak the compute-buffer allocator has to cover - n_copies times over under
pipeline parallelism.

Keep the growable array (fixes the original >30-input assert), and cut on
the constant again.

Cherry-picked from nathanw1014/llama.cpp strix-halo-vulkan (76ad2ba).
cubetitled-ui pushed a commit to cubetitled-ui/llama.cpp that referenced this pull request Aug 14, 2026
* ggml: use dynamic allocation for split graph inputs

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant

* cont : clean-up

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
brittlewis12 pushed a commit to brittlewis12/llama.cpp that referenced this pull request Aug 17, 2026
* ggml: use dynamic allocation for split graph inputs

Replace fixed-size GGML_SCHED_MAX_SPLIT_INPUTS arrays with dynamically
allocated buffers in the backend scheduler. This fixes crashes when
loading wide MoE models (Gemma 4, Qwen MoE, Mixtral, DeepSeek) on
multi-backend setups where graph splits exceed 30 input tensors.

- split->inputs: dynamic array with grow-on-demand
- sched->graph_inputs: dynamic array with grow-on-demand
- graph_size calculation now uses actual input count instead of fixed constant

* cont : clean-up

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ggml changes relating to the ggml tensor library for machine learning

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants