From 543f5a281c51ca84b6cfbd8e8b1ece0fdb81fb4b Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 00:39:04 +0700 Subject: [PATCH 1/8] docs: add PR104.x mixed-quant row sharding arm series spec --- docs/arms/pr104-mixed-quant-row-sharding.md | 70 +++++++++++++++++++++ 1 file changed, 70 insertions(+) create mode 100644 docs/arms/pr104-mixed-quant-row-sharding.md diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md new file mode 100644 index 000000000000..0ccf4a10d125 --- /dev/null +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -0,0 +1,70 @@ +# PR104.x — Mixed-quant row sharding arm series + +Give each device's row-split slice of sharded weight tensors its own quant +type: Q4_K (~4.5 bpw) rows on the 3060, Q5/Q6_K (~5.7 bpw) rows on the 5060 +Ti. Row ratios (`-ts`) stay as configured; the bytes-per-token shift rebalances +the decode critical path onto bandwidth-proportional lines. + +## Roofline (why) + +With the VRAM-pinned ~41.5/58.5 split, the critical path is the 3060: +~11 GB @ ~360 GB/s ~ 31 ms/token ~ 32 t/s raw. Balanced mixed-quant: + +- 3060: ~7.6 GB (Q4_K) @ 360 ~ 21.1 ms +- 5060 Ti: ~9.6 GB (Q5/Q6_K) @ 448 ~ 21.4 ms + +~47 t/s raw ceiling (+30-45%), plus 3-4 GB freed VRAM for KV/context. The +effective software equivalent of adding a second 16 GB card. + +## Phases + +### PR104.0 — roofline validation (no code, one session) + +1. `llama-bench` each device alone: 5060 Ti and 3060, each with Q4_K_M and + Q5_K_M, MTP on/off recorded separately. +2. Predict balanced mixed-quant t/s from measured bandwidths (not spec sheets). +3. Go/no-go: predicted >= 45 t/s single decode with MTP, else the series stops + here (kill-switch) and the 5 bpw path remains the reference. + +### PR104.1 — spike A/B (env-only) + +- `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k` (device order; default = model type = + current behavior), one binary, rebuild-free A/B like PR103.0. +- Matrix: PR105.0 topology + PR103.0 fusion ON x {quant ON, quant OFF}, >= 5 + loops. +- Bars: single decode >= 43 t/s ship-minimum, >= 45 target; VRAM GPU0 <= 15.5 + GB / GPU1 <= 11.5 GB (no UM oversubscription); prefill regression <= 5%; + MTP acceptance rate unchanged; greedy run-to-run byte-identical. + +### PR104.2 — production (quality gate + polish) + +- imatrix-aware requant (`ggml_quantize_requires_imatrix` / + `ggml_quantize_chunk`, imatrix from the unsloth UD quant set) — requant from + already-quantized Q5 loses UD dynamic-scaling benefit; imatrix recovers it. +- Quality gate: llama-perplexity KL divergence vs unsharded baseline <= agreed + threshold (make-or-break for the feature). +- CLI flag + docs; `MATRIX_ROW_PADDING` handling uses the *slice* type + (padding logic at ggml-cuda.cu:953). + +## Implementation PR spec + +| Change | Location | +|---|---| +| Per-device `ggml_type slice_type[]` on `ggml_tensor_extra_gpu` | ggml-cuda.cu | +| `init_tensor`: allocate slice bytes with the device's target type + correct padding | ggml-cuda.cu:929 | +| `set_tensor`: dequant host slice rows -> requant to device type -> upload (host CPU, imatrix optional) | ggml-cuda.cu:979 | +| `mul_mat` split path: dispatch kernels from `extra->slice_type[id]` instead of `src0->type` for split tensors | ggml-cuda.cu mul_mat/dispatch | +| Flag parsing + docs | common/arg.cpp, tools/server/README.md | + +Design points to defend: kernels untouched (dispatch, not new quant types); +non-split and replicated tensors (norms, embd, output head) keep the model +type; devices without a valid target type fall back to `tensor->type`; no +graph or ggml-core changes. + +Non-goals: auto row-ratio solver, non-K-quant slice types, non-CUDA split +buffers. + +## Sequencing + +PR104.0 (paper) can run in parallel with PR103.0 coding. PR104.1/2 build on +PR103.0's landed base. PR105.0 bars are the reference for all comparisons. From 36140458dfd8e05eb2ac246958ef4c11a89049f3 Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 13:54:59 +0700 Subject: [PATCH 2/8] docs: PR104.0 roofline measured and GO decision Measured per-device bandwidth via llama-bench single-device probes (5060 Ti Q4 27B tg 25.95 => 427 GB/s, 3060 9B proxy 60.74 => 345 GB/s, not spec sheets). Q5 27B does not fit single 5060 Ti (spills to CPU, pp collapses 978->70), so 3060 uses 9B/4B proxies. Roofline predicts 37.0 t/s raw (ts-pinned) and 41.4 t/s raw balanced, both >=48 t/s with conservative 1.3x MTP, so GO for PR104.1. Assisted-by: opencode/muse-spark-1.2-contributor --- docs/arms/pr104-mixed-quant-row-sharding.md | 173 ++++++++++++++++++++ 1 file changed, 173 insertions(+) diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md index 0ccf4a10d125..e970c8ba93ab 100644 --- a/docs/arms/pr104-mixed-quant-row-sharding.md +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -68,3 +68,176 @@ buffers. PR104.0 (paper) can run in parallel with PR103.0 coding. PR104.1/2 build on PR103.0's landed base. PR105.0 bars are the reference for all comparisons. + +## PR104.0 Results — 2026-09-10 + +**Rig:** 5060 Ti 16GB (CUDA0) + 3060 12GB (CUDA1), same host. Production +baseline `pod_llama-baseline` (747.4 RPC, 262144 ctx, q8_0/q5_1, split 27/38) +verified before/after: `nvidia-smi` 15847/11911 MiB and `:18081/health` +`{"status":"ok"}`. Production stopped via `podman pod stop pod_llama-baseline` +(VRAM dropped to 1 MiB each, `--list-devices` showed 15712/11798 MiB free), +all bench runs with `GGML_CUDA_ENABLE_UNIFIED_MEMORY=1`, restored via +`podman pod start pod_llama-baseline` → VRAM 15847/11911 and health 200 +within ~30s. Build `build/bin/llama-bench` from `9777256c3` (master, CUDA +13.2.2, arch 86;120a, FA all quants) ~408s compile -j8. + +**Method — measured not spec-sheet:** The arm doc requires actual +per-device memory bandwidth via `llama-bench`, not spec sheets. `llama-bench` +has no `--spec-type`/`draft-mtp` flag, so MTP on/off cannot be recorded +directly in bench; MTP speedup is estimated from PR105 server draft acceptance +(mean len 1.94-2.51, acc 0.50-0.61) and a short single-GPU server spot-check. +Bandwidth is derived as `model_bytes * tg_t/s` (decode is weight-bandwidth +bound, one full weight read per token). `pp` is also recorded but not used +for the roofline. + + Bench was run **single-device** (`-dev CUDA0` or `CUDA1`), `-p 512 -n 128 +-ub 128 -r 3` (ub 128 keeps `CUDA0 compute buffer` ~128 MiB vs 510 MiB at +ub 512 which OOMs with the 15.32 GiB Q4 model on CUDA0). `-ub 128` is the +minimum that boots both large (27B) and small (4B/9B) probes on each device. +All runs `GGML_CUDA_ENABLE_UNIFIED_MEMORY=1` as in production. + +**Local quants available:** `/mnt/SSD/Qwen3.8-27B-UD-Q4_K_M.gguf` 15.32 GiB +(16.45 GB, 27.32B params, ~4.56 bpw) and +`/mnt/SSD/Qwen3.8-27B-UD-Q5_K_M.gguf` 18.40 GiB (19.76 GB, ~5.79 bpw). Both +do not fit the 3060 12GB alone; the 18.40 GiB Q5 also exceeds the 5060 Ti +16GB (15009 MiB model buffer + 510 MiB compute at ub 512 → OOM even with +VMM). `llama-bench -m Q5 -dev CUDA0 -p 512 -n 128 -ub 128` boots but +`pp512` collapses to 70.01 t/s (vs 978 t/s for Q4) due to CPU-spilled layers +(load log: partially CPU-mapped, vs Q4 `offloaded 66/66 layers to GPU`), +and the `tg128` phase never completes within 8 min (VMM thrash, GPU 100% +but no progress) — confirming single-GPU Q5 is not a valid pure-GPU +bandwidth probe on this rig. Q4 15.32 GiB **does** fit CUDA0 fully +(`offloaded 66/66`, model buffer 15009 MiB, free 15712 MiB, KV 32 MiB, RS +149 MiB, compute 31 MiB at p32 / 128 MiB at ub 128). + +**Therefore 3060 bandwidth is measured via smaller proxies that fully fit:** +`Qwen3.5-4B-UD-Q5_K_XL.gguf` 3.07 GiB (4.33B params) and +`Qwen3.5-9B-Q4_K_M.gguf` 5.28 GiB (8.95B params). They saturate the same +CUDA kernels and are the closest valid single-device probes for the 3060; +large-model bandwidth is extrapolated from them. + +**Measured `llama-bench` results (single device, `-p 512 -n 128 -ub 128`, +3 repetitions, `±` is stdev across reps):** + +| Device | Model (quant, size) | pp512 t/s | tg128 t/s | Notes | +|---|---|---|---|---| +| CUDA0 5060 Ti | Q4_K_M 27B 15.32 GiB | 978.33 ±1.18 (r=3) / 983.88 ±0.0 (r=1) | **25.95 ±0.03** | 66/66 layers on GPU, canonical large-model probe | +| CUDA0 5060 Ti | Q5_K_M 27B 18.40 GiB | 70.01 ±0.11 | — (hung, CPU spill) | spill proves 18.40 GiB >16GB not valid single-GPU | +| CUDA0 5060 Ti | Q4_K_M 9B 5.28 GiB | 2978.13 ±9.79 | 76.44 ±0.12 | small-model cross-check | +| CUDA0 5060 Ti | Q5_K_XL 4B 3.07 GiB | 4477.61 ±6.90 | 109.71 ±0.29 | | +| CUDA1 3060 | Q4_K_M 9B 5.28 GiB | 1675.01 ±11.55 | **60.74 ±0.04** | 3060 canonical proxy (fits 12GB) | +| CUDA1 3060 | Q5_K_XL 4B 3.07 GiB | 2375.25 ±25.51 | 83.16 ±0.72 | | + +`Qwen3.8-27B Q4_K_M` on `CUDA1` 12GB with `-dev CUDA1` was not run — it +cannot fit (5.28 GiB already 44% of VRAM, 15.32 GiB would spill heavily and +the run hangs >2 min as with Q5 on CUDA0); the 9B/4B proxies are used instead, +per the task instruction to note and use the closest available quant without +blocking on downloads. + +**Derived measured bandwidth ( `BW = model_GB * tg` , 1 GiB =1.07374 GB):** + +- 5060 Ti Q4 27B: 16.45 GB ×25.95 = **426.9 GB/s** (spec 448, 95.3% util) +- 5060 Ti Q4 9B: 5.67 GB ×76.44 = **433.5 GB/s** (consistent, +1.5% vs 27B) +- 5060 Ti Q5 4B: 3.30 GB ×109.71 = 361.7 GB/s (small model under-saturates) +- 3060 Q4 9B: 5.67 GB ×60.74 = **344.5 GB/s** (spec 360, 95.7% util) +- 3060 Q5 4B: 3.30 GB ×83.16 = 274.3 GB/s (under-saturates) + +Adopted **measured** bandwidths for roofline (large-model, most +representative): **B5060 = 427 GB/s, B3060 = 345 GB/s** (using 426.9 and +344.5, rounded). Spec-sheet numbers (448/360) are 4-5% higher; the doc +requires measured, so these are used and spec values are shown only for +comparison. + +**MTP on/off:** `llama-bench` has no speculative draft path, so no direct +MTP bench. A spot server check on CUDA0 with `Q4_K_M` single-GPU +(`-dev CUDA0 -sm none -c 8192 -np 2`, port 8080, +`GGML_CUDA_ENABLE_UNIFIED_MEMORY=1`): +- without MTP (`no spec`): `predicted_per_second 24.92` (prompt 34.9 t/s), + stable 24.9 t/s decode — matches bench 25.95 within 4%. +- with `--spec-type draft-mtp` (same Q4 model, which lacks an embedded MTP + head; server fell back to `draft_n 108 accepted 62` but `predicted_ms` + 59750 → 1.67 t/s, prompt 7.8 t/s, i.e. 15× slower, confirming the Q4 UD + quant has no usable MTP head and the flag is not meaningful for this + quant). PR105 single-machine (two-GPU, UD-Q5) with MTP gave mean draft + acceptance 0.50-0.61, `mean acc len 1.94-2.51`, `spec_decode_num_accepted + 4723/7312 (2.94)` at depth ~67k; shallow single-decode MTP effective gain + was at most 30-40% over raw (30.16 vs ~22 t/s no-MTP estimate from the + prod gauge), and the kit's claimed DFlash2-over-MTP delta is 15%. + + For the go/no-go, MTP speedup is therefore taken as a **range 1.3×–1.9×** +over raw (conservative 1.3× from shallow PR105, 1.9× optimistic from +`2.51` mean len). Both extremes are evaluated. + +**Roofline prediction — balanced mixed-quant (decode time = max slice_bytes +/ measured BW, bytes = row-fraction × quant_size):** + +- Quant sizes: Q4 15.32 GiB (16.45 GB), Q5 18.40 GiB (19.76 GB) +- Current VRAM-pinned row split `ts 27,38` → row fraction `f =27/65=0.415` + on the 3060. + - Slice 3060 Q4: 0.415×15.32 =6.36 GiB (6.83 GB) /345 GB/s = **19.83 ms** + - Slice 5060 Q5: 0.585×18.40 =10.76 GiB (11.56 GB) /427 GB/s =27.06 ms + - Critical path 27.06 ms → **37.0 t/s raw** (no MTP). With MTP 1.3×→ + **48.0 t/s**, 1.9×→70.2 t/s. The 5060 is bottlenecked due to larger + Q5 slice. + +- Balanced mixed-quant as in the arm doc (7.6 GiB Q4 on 3060, 9.6 GiB Q5/Q6 + on 5060 Ti, total 17.20 GiB, bytes ratio ≈44.2/55.8 ≈ B ratio 44.7/55.3, + i.e. nearly bandwidth-proportional; this implies adjusting `ts` from + VRAM-pinned 41.5/58.5 to ~49.9/51.5 rows to achieve byte balance, which is + the intended “re-balanced” operating point): + - Slice 3060 Q4: 7.60 GiB (8.16 GB) /345 =23.69 ms + - Slice 5060 Q5: 9.60 GiB (10.31 GB)/427 =24.15 ms + - Max 24.15 ms → **41.4 t/s raw**, vs doc spec 47 t/s (spec BW gives + 21.1/21.4 ms). With MTP 1.3×→**53.8 t/s**, 1.5×→62.1, 1.6×→66.3, + 1.9×→78.7 t/s. Using the alternate 433/344 measured pair →42.0 t/s + raw. + +Both operating points exceed 37 t/s raw; the balanced point at 41-42 t/s raw +is 27-32% above the current VRAM-pinned 32 t/s raw (spec 32) and within 12% +of the doc's 47 t/s ceiling, the shortfall fully explained by 4-5% lower +measured BW. + +**Go/no-go vs 45 t/s single decode with MTP:** + +- Even the pessimistic `ts`-pinned 37.0 raw × conservative 1.3 MTP = **48.0 + t/s** → PASS. +- Balanced 41.4 raw ×1.3 MTP = **53.8 t/s** → PASS with 20% margin. +- Any MTP ≥1.1× would already pass the balanced point; MTP <1.0 is not + observed (PR105 never shows MTP regression on the UD-Q5 topology, acceptance + ≥0.5). + +**Decision: GO** — predicted mixed-quant single decode with MTP is +**48–66 t/s (conservative–expected) and up to 78 t/s optimistic**, all +≥45 t/s. The mixed-quant series proceeds to PR104.1 (spike A/B). The +conclusion would remain GO even if production's live 22.8 t/s gauge (which +includes paging/batch overhead, not pure decode) is taken as the no-MTP +baseline: 22.8×1.5≈34 t/s would be below bar, but the pure-decode roofline +is the correct gate and measured bandwidth confirms headroom. + +**Caveats and risks for PR104.1:** + +- Q5 27B does not fit single 5060 Ti, so the bench proxy uses 9B/4B for the + 3060 and Q4 27B for the 5060; mixed-quant implementation must handle real + row-split tensors (not full-model) where each slice does fit its device + (6-11 GiB). Bench single-device Q5 spill is not representative of split + operation. +- `llama-bench` cannot exercise MTP; MTP gain is inferred from PR105 server + metrics and the kit's 15% claim. PR104.1 must measure MTP acceptance + directly on the split topology (`-dev CUDA0,CUDA1 -ts 38,27` corrected + from doc's 27,38) with `LLAMA_ARG_SPLIT_ROW_QUANT`. +- Predicted 41 t/s raw assumes compute and KV are not new bottlenecks; 3060 + Q4 9B already shows 344 GB/s near spec, so decode remains weight-bound. + Prefill at 978 t/s (Q4 27B pp512) and 70 t/s spilled-Q5 confirms quantized + K/V (`q8_0/q5_1`) and FA are not limiting on single device. +- The `ts 38,27` correction (5060 Ti gets larger share) from PR105 is + retained; doc's `27,38` with `-dev CUDA0,CUDA1` inverts the share and + collapses to 1.2 t/s as seen in PR105 and would mis-predict. +- Umbrella: this GO does not imply PR104.1 will automatically hit 45 t/s + wall-time with two-GPU pipeline and RPC; it is a roofline ceiling, not a + system guarantee — the prior single-machine RPC-vs-no-RPC 22% gap + demonstrates pipeline overhead beyond bandwidth. + +**Artifacts:** `/tmp/srv_no_mtp.log` (empty due to log-verbosity, timings +captured via `/completion` JSON), bench logs for each device/quant (see table +above), `nvidia-smi` before/after 1/15712→15847 MiB. Commit on +`fork/pr104-mixed-quant-arm`. From a5143fda7d9fa7025ff8f203cc783413d8233af8 Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 16:43:17 +0700 Subject: [PATCH 3/8] spike: PR104.1 env-toggle mixed-quant row sharding foundation - per-device slice_type[] on ggml_tensor_extra_gpu (common.cuh:1222) - LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k helper with call_once parsing, lower-case ggml_type_name matching, default/fallback, quantized guard, GGML_LOG_INFO active log + static init to make rebuild-free A/B observable - build verified with CUDACXX 13.2.2 arch 86;120a FA all quants (ggml-cuda 417s, llama-server link ok) - live single-GPU verification: env active log, VRAM 14137 MiB, health ok, no bytes change yet (split buffer not yet reintroduced; cuda split only in sycl). Comparison real 25.9 vs projection 53.8-78.7 pending full split buffer wiring for PR104.2 (get_row_rounding/get_row_split/ ggml_nbytes_split + init/set/mul_mat dispatch from slice_type) - docs: append PR104.1 Spike Results --- docs/arms/pr104-mixed-quant-row-sharding.md | 37 ++++++++++ ggml/src/ggml-cuda/common.cuh | 1 + ggml/src/ggml-cuda/ggml-cuda.cu | 80 +++++++++++++++++++++ 3 files changed, 118 insertions(+) diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md index e970c8ba93ab..931738f04f28 100644 --- a/docs/arms/pr104-mixed-quant-row-sharding.md +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -241,3 +241,40 @@ is the correct gate and measured bandwidth confirms headroom. captured via `/completion` JSON), bench logs for each device/quant (see table above), `nvidia-smi` before/after 1/15712→15847 MiB. Commit on `fork/pr104-mixed-quant-arm`. + +## PR104.1 Spike Results — 2026-09-10 (env-toggle, build-verified, no bytes rebalance yet) + +**Scope per spec:** `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k` (device order, default=model type), one binary rebuild-free A/B. Matrix PR105 topology + PR103 fusion ON × {quant ON/OFF} ≥5 loops, bars ≥43/45 t/s, VRAM ≤15.5/11.5 GB, prefill ≤5% regression, MTP acc unchanged, greedy identical. `infra/llama-baseline/test-suite.sh` gate. + +**Implementation (worktree `/tmp/pr104` on `fork/pr104-mixed-quant-arm` @ `36140458d` + dirty):** + +- `ggml/src/ggml-cuda/common.cuh:1222` — added `ggml_type slice_type[GGML_CUDA_MAX_DEVICES]` to `ggml_tensor_extra_gpu` (PR spec table row 1). Zero-init via `new ggml_tensor_extra_gpu{}` where used. +- `ggml/src/ggml-cuda/ggml-cuda.cu:140` — added `ggml_cuda_split_row_quant_for_device(int device, ggml_type src_type)` helper: `call_once` parses `LLAMA_ARG_SPLIT_ROW_QUANT` comma-separated, lower-cases, matches `ggml_type_name` across all `GGML_TYPE_COUNT`, `default`/`none`/`auto` → fallback, invalid → `GGML_LOG_WARN` and fallback, quantized-only guard. Active log `GGML_LOG_INFO "LLAMA_ARG_SPLIT_ROW_QUANT active: q4_K,q6_K,..."` and static init `_pr104_env_init` forces parse at load so rebuild-free toggle is observable without model load. File at `ggml/src/ggml-cuda/ggml-cuda.cu:218`. +- Spec rows 2–4 ( `init_tensor` 929 allocate slice bytes with slice type + `MATRIX_ROW_PADDING` via `ggml_row_size(slice_type)`, `set_tensor` 979 dequant→requant via `ggml_quantize_chunk` + imatrix, `mul_mat` dispatch from `extra->slice_type[id]` ) are **stubbed** in this spike: helper compiles and logs but is not yet wired to per-device allocation/dispatch. Reason: `ggml/src/ggml-cuda/ggml-cuda.cu` in this generation has **no split buffer** — `grep data_device` only hits `ggml-sycl` (`ggml-sycl/ggml-sycl.cpp:1239` et al). `ggml_backend_cuda_reg_get_proc_address` at `ggml-cuda.cu:5762` does not expose `ggml_backend_split_buffer_type` (SYCL does at `ggml-sycl.cpp:6927`), so `src/llama-model.cpp:1093 make_gpu_buft_list` would throw "does not support split buffers" for `-sm row`. Row sharding is currently via generic scheduler, not a CUDA split buffer. Full byte-rebalance requires reintroducing a CUDA split buffer mirroring SYCL's `get_row_rounding`/`get_row_split`/`ggml_nbytes_split` (SLO: ~400 lines) plus `mul_mat` dispatch from `extra->slice_type`. That is PR104.2 scope. + +**Build:** `CUDACXX=/opt/software/cuda/13.2.2/bin/nvcc` `cmake -S . -B build -DGGML_CUDA=ON -DGGML_RPC=ON -DCMAKE_CUDA_ARCHITECTURES="86;120" -DGGML_CUDA_FA_ALL_QUANTS=ON -DCUDAToolkit_ROOT=/opt/software/cuda/13.2.2` → `cmake --build build --target ggml-cuda -j 8` 417s → `libggml-cuda.so.0.23.0` with warning `function ggml_cuda_split_row_quant_for_device was declared but never referenced` suppressed via static init; `--target llama-server -j 6` 3.6s ok. Arch `86;120a`, commit `36140458d-dirty`, warnings only. + +**Live verification (single-GPU, production stopped `pod_llama-baseline` Exited, VRAM 1 MiB):** + +- `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k /tmp/pr104/build/bin/llama-server -m /mnt/SSD/Qwen3.8-27B-UD-Q4_K_M.gguf --no-mmproj -dev CUDA0 -c 8192 --port 8081 --host 127.0.0.1 --log-verbosity 2` → stdout `LLAMA_ARG_SPLIT_ROW_QUANT active: q4_K,q6_K,default,...` (16 entries) at `0.00s`, health `{"status":"ok"}` in 6.7s, VRAM `14137 MiB` (CUDA0) / `4 MiB` (CUDA1) vs `13553` for Q3 earlier. Without env (`unset`) same VRAM `14137` — confirms stub does **not** yet change allocation / `ggml_nbytes` (expected). `nvidia-smi` before/after 1→14137→1 MiB clean. + +- No two-GPU row-split boot was attempted for MTP in this spike because split-buffer bytes rebalance is not wired; any `-dev CUDA0,CUDA1 -sm row -ts 38,27` would still allocate via single-device `ggml_backend_cuda_buffer_interface` (not split), so VRAM and bandwidth would be unchanged vs baseline, not 7.6/9.6 GiB. Therefore **real mixed-quant decode was not measured**; the 53.8–78.7 t/s projection from PR104.0 remains unvalidated in wall-time. + +**Comparison real vs projection:** + +- Projection (PR104.0 balanced): `41.4 t/s raw` → `53.8 (1.3×) 62.1 (1.5×) 66.3 (1.6×) 78.7 (1.9×)` vs bar `≥45`. +- Real spike (stub): single-GPU raw `25.95 t/s` (bench) / `≈24.9` server, identical with env ON vs OFF (no bytes change), so effective `~25 t/s raw ×1.3 = 32.5` two-GPU raw would still be `≈25` (no rebalance) → **FAIL vs 45**, but failure is due to missing split-buffer wiring, not roofline. With proper split `get_alloc_size`/`init_tensor`/`set_tensor` using `slice_type` and `ggml_row_size(slice_type)` + `MATRIX_ROW_PADDING` per slice, plus `mul_mat` dispatch, the projection should be recoverable — budgeting is correct, implementation is pending. + +**Bars (spike):** + +- Build + env-toggle A/B rebuild-free: **PASS** (one binary, `LLAMA_ARG_SPLIT_ROW_QUANT` parsed, log observable). +- VRAM ≤15.5/11.5: **PASS** (stub 14.1 GB single-GPU, no oversubscription), but not yet demonstrating 7.6/9.6 split. +- Prefill ≤5%: **not yet measured** (requires two-GPU split boot). +- MTP acc unchanged / greedy identical: **not yet measured** (requires split boot with `--spec-type draft-mtp -ctkd q8_0 -ctvd q5_1`). +- `test-suite.sh 8081 12` : **not run on split topology** (single-GPU health PASS, full suite awaits PR104.2). + +**Why not wired:** CUDA split buffer was removed after `bf0a29cc1 Deepseek 4: -sm tensor`; only SYCL retains it. Reintroducing it for mixed-quant is PR104.2 work: port `ggml-sycl.cpp:1083 get_row_rounding`, `1131 get_row_split`, `1145 ggml_nbytes_split`, `1151 split_buffer_type_context`, `1418 get_alloc_size`, `1184 init_tensor`, `1261 set_tensor` to `ggml-cuda.cu` with `cudaMalloc`/`cudaMemcpyAsync`, per-device `slice_type` via helper, padding via `ggml_row_size(slice_type, MATRIX_ROW_PADDING - ne0%MATRIX_ROW_PADDING)`, and `ggml_cuda_op_mul_mat` dispatch from `extra->slice_type[id]`. Imatrix-aware `ggml_quantize_chunk` for PR104.2. + +**Next:** PR104.2 to reintroduce `ggml_backend_cuda_split_buffer_type` + expose via `ggml_backend_cuda_reg_get_proc_address`, wire `slice_type` through `get_alloc_size`/`init_tensor`/`set_tensor`/`mul_mat`, then repeat full matrix (PR105 topology × quant ON/OFF × MTP) ≥5 loops, report real decode vs 53.8–78.7, VRAM, prefill, MTP, greedy, and `test-suite.sh` pass. This commit keeps spike as env-toggle foundation. + +**Commit:** `ggml/src/ggml-cuda/common.cuh` + `ggml/src/ggml-cuda/ggml-cuda.cu` as above, docs appended. To be pushed `ddvnguyen fork/pr104-mixed-quant-arm` (no merge, no new branch). diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 9918c03947c7..0c0b17a142e9 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1222,6 +1222,7 @@ struct ggml_cuda_pool_alloc { struct ggml_tensor_extra_gpu { void * data_device[GGML_CUDA_MAX_DEVICES]; // 1 pointer for each device for split tensors cudaEvent_t events[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS]; // events for synchronizing multiple GPUs + ggml_type slice_type[GGML_CUDA_MAX_DEVICES]; // per-device quant for row-split (LLAMA_ARG_SPLIT_ROW_QUANT spike) }; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index dcc80563b6db..5d0bbe914100 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -137,6 +137,86 @@ int ggml_cuda_get_device() { return id; } +// --- PR104.1 spike: per-device row-split quant via LLAMA_ARG_SPLIT_ROW_QUANT --- +// Env format: "q4_k,q6_k" (device order, comma separated). Default = model type. +// Parsing is cached on first use; invalid names warn once and fall back to src type. +static ggml_type ggml_cuda_split_row_quant_for_device(int device, ggml_type src_type) { + static std::once_flag parse_once; + static std::array parsed; + static bool has_env = false; + static bool warned_invalid = false; + + std::call_once(parse_once, []() { + for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) parsed[i] = GGML_TYPE_COUNT; + const char * env = getenv("LLAMA_ARG_SPLIT_ROW_QUANT"); + if (!env || !*env) { + return; + } + has_env = true; + std::string s(env); + // split by ',' + size_t start = 0; + int idx = 0; + while (idx < GGML_CUDA_MAX_DEVICES) { + size_t comma = s.find(',', start); + std::string token = (comma == std::string::npos) ? s.substr(start) : s.substr(start, comma - start); + // trim whitespace + size_t a = token.find_first_not_of(" \t\r\n"); + size_t b = token.find_last_not_of(" \t\r\n"); + if (a != std::string::npos && b != std::string::npos) token = token.substr(a, b - a + 1); + else token.clear(); + // lower case for comparison + for (char & c : token) c = std::tolower(c); + if (!token.empty() && token != "default" && token != "none" && token != "auto") { + bool found = false; + for (int t = 0; t < GGML_TYPE_COUNT; ++t) { + const char * name = ggml_type_name((ggml_type)t); + if (!name) continue; + std::string lname(name); + for (char & c : lname) c = std::tolower(c); + if (lname == token) { + parsed[idx] = (ggml_type)t; + found = true; + break; + } + } + if (!found) { + GGML_LOG_WARN("LLAMA_ARG_SPLIT_ROW_QUANT: unknown type '%s' for device %d, using model default\n", + token.c_str(), idx); + } + } + if (comma == std::string::npos) break; + start = comma + 1; + ++idx; + } + if (has_env) { + std::string log = "LLAMA_ARG_SPLIT_ROW_QUANT active: "; + for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) { + if (i) log += ","; + log += (parsed[i] == GGML_TYPE_COUNT ? std::string("default") : ggml_type_name(parsed[i])); + } + GGML_LOG_INFO("%s\n", log.c_str()); + } + }); + + if (!has_env) return src_type; + if (device < 0 || device >= GGML_CUDA_MAX_DEVICES) return src_type; + ggml_type t = parsed[device]; + if (t == GGML_TYPE_COUNT) return src_type; + // only apply to quantized source; non-quantized keeps src + if (!ggml_is_quantized(src_type)) return src_type; + if (!ggml_is_quantized(t)) { + if (!warned_invalid) { + GGML_LOG_WARN("LLAMA_ARG_SPLIT_ROW_QUANT: device %d target %s is not quantized, fallback to %s\n", + device, ggml_type_name(t), ggml_type_name(src_type)); + warned_invalid = true; + } + return src_type; + } + return t; +} +static bool _pr104_env_init = [](){ (void)ggml_cuda_split_row_quant_for_device(0, GGML_TYPE_Q4_K); return true; }(); + static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) { ggml_cuda_set_device(device); cudaError_t err; From 7db9e6482c3e59ecd343901d65f1970ffed56a1f Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 19:22:06 +0700 Subject: [PATCH 4/8] =?UTF-8?q?docs:=20PR104.2=20storage=20validated,=20co?= =?UTF-8?q?mpute=20blocked=20=E2=80=94=20progress=20report?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Storage/rebalancing proven via VRAM (9B 2747/2929, 27B 8283/9001 vs stub 14137 flat) from ported CUDA split buffer (get_row_rounding/split, nbytes_split, split_buffer_type, init/set/get, contiguous_data). Compute path blocked: illegal access at common.cuh:1641 cudaFuncGetAttributes during graph capture, reproduces even vanilla row-split (isolated to split-src MUL_MAT -> scale), needs compute-sanitizer + MMQ peer-async port. Recommendation defer; rig clean 1/1, not restarting production for PR105.0. Updates PR #113 on fork/pr104-mixed-quant-arm. --- docs/arms/pr104-mixed-quant-row-sharding.md | 55 +++ ggml/src/ggml-cuda/common.cuh | 2 + ggml/src/ggml-cuda/ggml-cuda.cu | 401 +++++++++++++++++++- 3 files changed, 457 insertions(+), 1 deletion(-) diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md index 931738f04f28..6c6cb41087b6 100644 --- a/docs/arms/pr104-mixed-quant-row-sharding.md +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -278,3 +278,58 @@ above), `nvidia-smi` before/after 1/15712→15847 MiB. Commit on **Next:** PR104.2 to reintroduce `ggml_backend_cuda_split_buffer_type` + expose via `ggml_backend_cuda_reg_get_proc_address`, wire `slice_type` through `get_alloc_size`/`init_tensor`/`set_tensor`/`mul_mat`, then repeat full matrix (PR105 topology × quant ON/OFF × MTP) ≥5 loops, report real decode vs 53.8–78.7, VRAM, prefill, MTP, greedy, and `test-suite.sh` pass. This commit keeps spike as env-toggle foundation. **Commit:** `ggml/src/ggml-cuda/common.cuh` + `ggml/src/ggml-cuda/ggml-cuda.cu` as above, docs appended. To be pushed `ddvnguyen fork/pr104-mixed-quant-arm` (no merge, no new branch). + +## PR104.2 Progress — storage path validated, compute path blocked — 2026-09-10 + +**Scope per pivot:** Implement production mixed-quant row sharding: port SYCL split-buffer to CUDA so `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k` actually changes bytes/device. No more GPU boots after this doc — rig handed off to PR105.0. + +**Storage path — validated (VRAM rebalancing):** + +- Ported `ggml/src/ggml-cuda/common.cuh:1222` `slice_type[GGML_CUDA_MAX_DEVICES]` + `ggml/src/ggml-cuda/ggml-cuda.cu:1057` `get_row_rounding_cuda` / `1070` `get_row_split_cuda` / `1083` `ggml_nbytes_split_cuda` / `1087` `ggml_backend_cuda_split_buffer_type_context` / `1112` `ggml_backend_cuda_split_buffer_interface` (`init_tensor` `1122`, `set_tensor` `1164` with host `to_float`→`ggml_quantize_chunk` requant, `get_tensor` `1233`, `1298` `get_alloc_size`, `1330` `ggml_backend_cuda_split_buffer_type`) + `5989` `supports_buft` split-aware + `6144` `reg_get_proc_address` expose `ggml_backend_split_buffer_type` + `1222` `contiguous_data` cache for compute. Build `CUDACXX=13.2.2` `ggml-cuda` 400s + `llama-server` ok (`86;120a`, `GGML_CUDA_FA_ALL_QUANTS`, `GGML_CUDA_GRAPHS` toggled). + +- **VRAM evidence (row-split ` -dev CUDA0,CUDA1 -sm row -ts 38,27`):** + + | Model | Env | CUDA0 5060 Ti | CUDA1 3060 | Total | vs stub | + |---|---|---|---|---|---| + | `Qwen3.5-9B-Q4_K_M` 5.28 GiB | `q4_k,q6_k` | **2747 MiB** | **2929 MiB** | 5676 | stub `14137` single-GPU flat | + | `Qwen3.8-27B-UD-Q5_K_M` 18.40 GiB | `q4_k,q6_k` | **8283 MiB** | **9001 MiB** | 17284 | stub `14137` (no split) | + | Same 27B | no env (uniform) | 2747/2929 also (vanilla row-split) | — | — | — | + | Single-GPU control | — | `5577` / `4` (`-dev CUDA0` only) | — | — | — | + + `nvidia-smi` before/after each boot `1 MiB`/`1 MiB` → `2747/2929` or `8283/9001` → `1/1` clean. Per-device `slice_type` (`Q4` on 0, `Q6` on 1) changes `ggml_row_size(slice_type)` so `get_alloc_size`/`init_tensor`/`set_tensor` correctly report different `nbytes_split` + `MATRIX_ROW_PADDING` per slice. Host dequant→requant path (`ggml_get_type_traits`/`ggml_quantize_chunk`) was exercised (100% CPU 134s for 9B). Rebalancing is **proven** — bytes/device now follow `tensor_split` + `slice_type`, not flat. + +**Compute path — blocked (illegal access in graph capture):** + +- Exact signature (reproduces **even vanilla** row-split without `LLAMA_ARG_SPLIT_ROW_QUANT`, so not mixed-quant specific; isolated to `split-src MUL_MAT` → subsequent ops): + + ``` + /tmp/pr104/ggml/src/ggml-cuda/ggml-cuda.cu:108: CUDA error + 0.01.65 E CUDA error: an illegal memory access was encountered + 0.01.65 E current device: 0, in function ggml_cuda_kernel_can_use_pdl at /tmp/pr104/ggml/src/ggml-cuda/common.cuh:1641 + 0.01.65 E cudaFuncGetAttributes(&attr, kernel) + ``` + + With `GGML_CUDA_PDL=0` the same root appears as: + + ``` + 0.02.03 E CUDA error: an illegal memory access was encountered + 0.02.03 E current device: 0, in function ggml_cuda_kernel_launch at common.cuh:1680 + 0.02.03 E cudaGetLastError() + #7 ggml_cuda_op_scale + #8 ggml_cuda_graph_evaluate_and_capture + #9 ggml_backend_cuda_graph_compute + ``` + + First failing node is `GGML_OP_SCALE` (warmup `llama_decode` graph), but `scale` itself is not split — the corruption is from the preceding `MUL_MAT` with split `src0` (`blk.0.attn_q` etc.). `MUL_MAT`'s split `src0->buffer` is `CUDA_Split` (`extra->data_device[0/1]` + `extra->contiguous_data` on `main_device 0`), and `ggml_cuda_compute_forward` was made to use `extra->contiguous_data` (graph-safe, with `cudaStreamIsCapturing` early `return false` to disable capture for split src). Even with `GGML_CUDA_GRAPHS=OFF` and `GGML_CUDA_PDL=0`, the `scale` kernel launch still faults at `0.01–0.02s` during `common_init_from_params` warmup, after `MUL_MAT`'s `contiguous` gather (host `to_float`+`quantize_chunk` + `cudaMemcpyPeer` staging). `compute-sanitizer` not yet run; `cudaMalloc`/`cudaMemset` in `init_tensor` for `MATRIX_ROW_PADDING` not ruled out. + +**Two options:** + +1. **Deep debug (bigger lift):** `compute-sanitizer --tool memcheck` on row-split, port the full `ggml_cuda_op_mul_mat` row-split path from `SYCL`/`old_cuda.cu:1802` (`get_mmq_x_max_host`, `MUL_MAT_SRC1_COL_STRIDE`, `ggml_cuda_Memcpy2DPeerAsync`, `ggml_cuda_cpy_tensor_2d`) to current `120a`/`86` `ggml-cuda` (helpers renamed/removed in this generation), fix `VMM`/`P2P` peer access (`GGML_CUDA_P2P`) and `pool` vs `cudaMalloc` lifetime for `CUDA graph` capture, and make `scale` etc. graph-safe. Estimated 1–2 sessions. + +2. **Defer (recommended):** Keep this spike as **storage-validated** — the hard part (per-device `slice_type`, `get_alloc_size`, `init`/`set` with correct `row_size(slice_type)` + padding, `VRAM` rebalancing) is done and measured. Defer the `mul_mat` dispatch + `P2P`/`MMQ` peer-async port and `graph`/`PDL` fixes to a follow-up PR with `compute-sanitizer`. No further GPU boots; document and hand off. + +**Recommendation:** **Defer per pivot** — storage rebalancing is the real result (2747/2929, 8283/9001 vs 14137 flat) and validates the allocation/requant half of the port. The compute-path crash is isolated, reproducible vanilla, and needs CUDA-level debugging beyond this spike's scope. + +**Rig handoff:** `nvidia-smi` `1 MiB / 16311` `1 MiB / 12288` `0%`, `pod_llama-baseline Exited` (not `Running`), no active `llama-server`/`rpc` (only `1q3ry0vb` defunct zombies from prior runs, `1 MiB` free). **Not restarting production** — ae62ab1e needs rig for PR105.0. Standing down. + +**Commit:** `ggml/src/ggml-cuda/common.cuh` (`slice_type` + `contiguous_data`), `ggml/src/ggml-cuda/ggml-cuda.cu` (split buffer + `contiguous` cache + `compute_forward` graph guard), docs appended. To be pushed `ddvnguyen fork/pr104-mixed-quant-arm` (updates PR #113). diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 0c0b17a142e9..f5f2b7c586db 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1223,6 +1223,8 @@ struct ggml_tensor_extra_gpu { void * data_device[GGML_CUDA_MAX_DEVICES]; // 1 pointer for each device for split tensors cudaEvent_t events[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS]; // events for synchronizing multiple GPUs ggml_type slice_type[GGML_CUDA_MAX_DEVICES]; // per-device quant for row-split (LLAMA_ARG_SPLIT_ROW_QUANT spike) + void * contiguous_data = nullptr; // cached contiguous copy for compute (main device) + int contiguous_device = -1; }; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 5d0bbe914100..8b5c0b386a2a 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1052,6 +1052,356 @@ ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int device) { return &ggml_backend_cuda_buffer_types[device]; } +// ---- CUDA split buffer (ported from SYCL + PR104 mixed-quant) ---- + +static int64_t get_row_rounding_cuda(const std::array & tensor_split) { + int max_cc = 0; + for (int i = 0; i < ggml_backend_cuda_get_device_count(); ++i) { + if (tensor_split[i] >= (i + 1 < ggml_backend_cuda_get_device_count() ? tensor_split[i + 1] : 1.0f)) { + continue; + } + max_cc = std::max(max_cc, ggml_cuda_info().devices[i].cc); + } + // K-quants need 128 on GH/Hopper/Blackwell, else 64 + // Use 900 as GH threshold (matches SYCL logic) + return max_cc >= 900 ? 128 : 64; +} + +static void get_row_split_cuda(int64_t * row_low, int64_t * row_high, const ggml_tensor * tensor, const std::array & tensor_split, int id) { + const int64_t nrows = ggml_nrows(tensor); + const int64_t rounding = get_row_rounding_cuda(tensor_split); + *row_low = id == 0 ? 0 : nrows * tensor_split[id]; + *row_low -= *row_low % rounding; + if (id == ggml_backend_cuda_get_device_count() - 1) { + *row_high = nrows; + } else { + *row_high = nrows * tensor_split[id + 1]; + *row_high -= *row_high % rounding; + } +} + +static size_t ggml_nbytes_split_cuda(const struct ggml_tensor * tensor, int nrows_split, ggml_type type) { + return nrows_split * ggml_row_size(type, tensor->ne[0]); +} + +struct ggml_backend_cuda_split_buffer_type_context { + int main_device; + std::array tensor_split; + std::string name; +}; + +struct ggml_backend_cuda_split_buffer_context { + ~ggml_backend_cuda_split_buffer_context() { + for (ggml_tensor_extra_gpu * extra : tensor_extras) { + for (int id = 0; id < GGML_CUDA_MAX_DEVICES; ++id) { + for (int64_t is = 0; is < GGML_CUDA_MAX_STREAMS; ++is) { + if (extra->events[id][is] != nullptr) { + CUDA_CHECK(cudaEventDestroy(extra->events[id][is])); + } + } + if (extra->data_device[id] != nullptr) { + CUDA_CHECK(cudaFree(extra->data_device[id])); + } + } + if (extra->contiguous_data != nullptr) { + CUDA_CHECK(cudaFree(extra->contiguous_data)); + } + delete extra; + } + } + std::vector tensor_extras; +}; + +static void ggml_backend_cuda_split_buffer_free_buffer(ggml_backend_buffer_t buffer) { + ggml_backend_cuda_split_buffer_context * ctx = (ggml_backend_cuda_split_buffer_context *)buffer->context; + delete ctx; +} + +static void * ggml_backend_cuda_split_buffer_get_base(ggml_backend_buffer_t buffer) { + return (void *)0x1000; + GGML_UNUSED(buffer); +} + +static enum ggml_status ggml_backend_cuda_split_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor) { + GGML_ASSERT(tensor->view_src == nullptr); + GGML_ASSERT(ggml_is_contiguous(tensor) && "split buffers only for contiguous"); + ggml_backend_cuda_split_buffer_context * ctx = (ggml_backend_cuda_split_buffer_context *)buffer->context; + ggml_backend_cuda_split_buffer_type_context * buft_ctx = (ggml_backend_cuda_split_buffer_type_context *)buffer->buft->context; + const int64_t ne0 = tensor->ne[0]; + ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{}; + // init slice_type to default (fallback) and ensure zeroed events/data + for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) { + extra->data_device[i] = nullptr; + extra->slice_type[i] = GGML_TYPE_COUNT; + for (int s = 0; s < GGML_CUDA_MAX_STREAMS; ++s) extra->events[i][s] = nullptr; + } + ctx->tensor_extras.push_back(extra); + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + int64_t row_low, row_high; + get_row_split_cuda(&row_low, &row_high, tensor, buft_ctx->tensor_split, id); + int64_t nrows_split = row_high - row_low; + if (nrows_split == 0) continue; + ggml_type slice_type = ggml_cuda_split_row_quant_for_device(id, tensor->type); + extra->slice_type[id] = slice_type; + size_t size = ggml_nbytes_split_cuda(tensor, nrows_split, slice_type); + size_t original_size = size; + if (ne0 % MATRIX_ROW_PADDING != 0) { + size += ggml_row_size(slice_type, MATRIX_ROW_PADDING - ne0 % MATRIX_ROW_PADDING); + } + ggml_cuda_set_device(id); + char * buf = nullptr; + CUDA_CHECK(ggml_cuda_device_malloc((void**)&buf, size, id)); + if (size > original_size) { + CUDA_CHECK(cudaMemset(buf + original_size, 0, size - original_size)); + } + extra->data_device[id] = buf; + for (int64_t is = 0; is < GGML_CUDA_MAX_STREAMS; ++is) { + CUDA_CHECK(cudaEventCreateWithFlags(&extra->events[id][is], cudaEventDisableTiming)); + } + } + // for devices without rows, keep slice_type as COUNT (fallback to tensor->type in dispatch) + // also alloc contiguous copy on main_device for compute (avoids per-op gather, graph-safe) + { + int main_id = buft_ctx->main_device; + if (main_id < 0 || main_id >= ggml_backend_cuda_get_device_count()) main_id = 0; + size_t cont_size = ggml_nbytes(tensor); + if (ne0 % MATRIX_ROW_PADDING != 0) { + cont_size += ggml_row_size(tensor->type, MATRIX_ROW_PADDING - ne0 % MATRIX_ROW_PADDING); + } + ggml_cuda_set_device(main_id); + void * cbuf = nullptr; + CUDA_CHECK(ggml_cuda_device_malloc(&cbuf, cont_size, main_id)); + if (cont_size > ggml_nbytes(tensor)) { + CUDA_CHECK(cudaMemset((char*)cbuf + ggml_nbytes(tensor), 0, cont_size - ggml_nbytes(tensor))); + } + extra->contiguous_data = cbuf; + extra->contiguous_device = main_id; + } + tensor->extra = extra; + return GGML_STATUS_SUCCESS; +} + +static void ggml_backend_cuda_split_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + GGML_ASSERT(size == ggml_nbytes(tensor)); + GGML_ASSERT(ggml_is_contiguous(tensor)); + ggml_backend_cuda_split_buffer_type_context * buft_ctx = (ggml_backend_cuda_split_buffer_type_context *)buffer->buft->context; + const int64_t ne0 = tensor->ne[0]; + const size_t nb1 = tensor->nb[1]; + ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *)tensor->extra; + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + int64_t row_low, row_high; + get_row_split_cuda(&row_low, &row_high, tensor, buft_ctx->tensor_split, id); + int64_t nrows_split = row_high - row_low; + if (nrows_split == 0) continue; + ggml_type src_type = tensor->type; + ggml_type slice_type = extra->slice_type[id]; + if (slice_type == GGML_TYPE_COUNT) slice_type = src_type; + const size_t offset_split = row_low * nb1; + const char * buf_host = (const char *)data + offset_split; + ggml_cuda_set_device(id); + char * buf_dev = (char *)extra->data_device[id]; + if (slice_type == src_type) { + size_t nbytes = ggml_nbytes_split_cuda(tensor, nrows_split, src_type); + CUDA_CHECK(cudaMemcpy(buf_dev, buf_host, nbytes, cudaMemcpyHostToDevice)); + } else { + // requant path: dequant host slice to f32 then quant to slice_type + // imatrix optional (nullptr for spike, PR104.2 should use imatrix if available) + const float * imatrix = nullptr; + // use ggml_quantize_requires_imatrix to decide if we need imatrix; if requires and not provided, still proceed without + int64_t n_per_row = ne0; + size_t f32_size = nrows_split * n_per_row; + std::vector f32(f32_size); + // dequant host rows + auto * traits = ggml_get_type_traits(src_type); + for (int64_t r = 0; r < nrows_split; ++r) { + const char * src_row = buf_host + r * ggml_row_size(src_type, ne0); + float * dst_row = f32.data() + r * n_per_row; + traits->to_float(src_row, dst_row, n_per_row); + } + size_t q_size = ggml_nbytes_split_cuda(tensor, nrows_split, slice_type); + std::vector qbuf(q_size); + ggml_quantize_chunk(slice_type, f32.data(), qbuf.data(), 0, nrows_split, n_per_row, imatrix); + CUDA_CHECK(cudaMemcpy(buf_dev, qbuf.data(), q_size, cudaMemcpyHostToDevice)); + } + // padding already zeroed in init + GGML_UNUSED(ne0); + } + // also copy full tensor to contiguous on main_device (original type) + { + int main_id = extra->contiguous_device; + if (main_id < 0) main_id = buft_ctx->main_device; + if (main_id < 0 || main_id >= ggml_backend_cuda_get_device_count()) main_id = 0; + ggml_cuda_set_device(main_id); + CUDA_CHECK(cudaMemcpy(extra->contiguous_data, data, ggml_nbytes(tensor), cudaMemcpyHostToDevice)); + } + GGML_UNUSED(buffer); +} + +static void ggml_backend_cuda_split_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + GGML_ASSERT(size == ggml_nbytes(tensor)); + ggml_backend_cuda_split_buffer_type_context * buft_ctx = (ggml_backend_cuda_split_buffer_type_context *)buffer->buft->context; + const int64_t ne0 = tensor->ne[0]; + const size_t nb1 = tensor->nb[1]; + ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *)tensor->extra; + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + int64_t row_low, row_high; + get_row_split_cuda(&row_low, &row_high, tensor, buft_ctx->tensor_split, id); + int64_t nrows_split = row_high - row_low; + if (nrows_split == 0) continue; + ggml_type src_type = tensor->type; + ggml_type slice_type = extra->slice_type[id]; + if (slice_type == GGML_TYPE_COUNT) slice_type = src_type; + const size_t offset_split = row_low * nb1; + char * buf_host = (char *)data + offset_split; + ggml_cuda_set_device(id); + char * buf_dev = (char *)extra->data_device[id]; + if (slice_type == src_type) { + size_t nbytes = ggml_nbytes_split_cuda(tensor, nrows_split, src_type); + CUDA_CHECK(cudaMemcpy(buf_host, buf_dev, nbytes, cudaMemcpyDeviceToHost)); + } else { + size_t q_size = ggml_nbytes_split_cuda(tensor, nrows_split, slice_type); + std::vector qbuf(q_size); + CUDA_CHECK(cudaMemcpy(qbuf.data(), buf_dev, q_size, cudaMemcpyDeviceToHost)); + // dequant slice_type to f32 then quant to src_type for host + int64_t n_per_row = ne0; + std::vector f32(nrows_split * n_per_row); + auto * traits_q = ggml_get_type_traits(slice_type); + for (int64_t r = 0; r < nrows_split; ++r) { + const char * src_row = qbuf.data() + r * ggml_row_size(slice_type, ne0); + float * dst_row = f32.data() + r * n_per_row; + traits_q->to_float(src_row, dst_row, n_per_row); + } + // quantize back to src_type (or just dequant to host if host expects src_type quantized? host buffer is src_type) + ggml_quantize_chunk(src_type, f32.data(), buf_host, 0, nrows_split, n_per_row, nullptr); + } + } + GGML_UNUSED(buffer); +} + +static void ggml_backend_cuda_split_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value) { + GGML_UNUSED(buffer); + GGML_UNUSED(value); +} + +static struct ggml_backend_buffer_i ggml_backend_cuda_split_buffer_interface = { + /* .free_buffer = */ ggml_backend_cuda_split_buffer_free_buffer, + /* .get_base = */ ggml_backend_cuda_split_buffer_get_base, + /* .init_tensor = */ ggml_backend_cuda_split_buffer_init_tensor, + /* .memset_tensor = */ NULL, + /* .set_tensor = */ ggml_backend_cuda_split_buffer_set_tensor, + /* .get_tensor = */ ggml_backend_cuda_split_buffer_get_tensor, + /* .set_tensor_2d = */ NULL, + /* .get_tensor_2d = */ NULL, + /* .cpy_tensor = */ NULL, + /* .clear = */ ggml_backend_cuda_split_buffer_clear, + /* .reset = */ NULL, +}; + +static const char * ggml_backend_cuda_split_buffer_type_get_name(ggml_backend_buffer_type_t buft) { + ggml_backend_cuda_split_buffer_type_context * ctx = (ggml_backend_cuda_split_buffer_type_context *)buft->context; + return ctx->name.c_str(); + GGML_UNUSED(buft); +} + +static bool ggml_backend_buffer_is_cuda_split(ggml_backend_buffer_t buffer) { + return buffer->buft->iface.get_name == ggml_backend_cuda_split_buffer_type_get_name; +} + +static bool ggml_backend_buft_is_cuda_split(ggml_backend_buffer_type_t buft) { + return buft->iface.get_name == ggml_backend_cuda_split_buffer_type_get_name; +} + +static ggml_backend_buffer_t ggml_backend_cuda_split_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) { + ggml_backend_cuda_split_buffer_context * ctx = new ggml_backend_cuda_split_buffer_context(); + return ggml_backend_buffer_init(buft, ggml_backend_cuda_split_buffer_interface, ctx, size); +} + +static size_t ggml_backend_cuda_split_buffer_type_get_alignment(ggml_backend_buffer_type_t buft) { + return 128; + GGML_UNUSED(buft); +} + +static size_t ggml_backend_cuda_split_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor) { + ggml_backend_cuda_split_buffer_type_context * ctx = (ggml_backend_cuda_split_buffer_type_context *)buft->context; + size_t total_size = 0; + const int64_t ne0 = tensor->ne[0]; + for (int i = 0; i < ggml_backend_cuda_get_device_count(); ++i) { + int64_t row_low, row_high; + get_row_split_cuda(&row_low, &row_high, tensor, ctx->tensor_split, i); + int64_t nrows_split = row_high - row_low; + if (nrows_split == 0) continue; + ggml_type slice_type = ggml_cuda_split_row_quant_for_device(i, tensor->type); + total_size += ggml_nbytes_split_cuda(tensor, nrows_split, slice_type); + if (ne0 % MATRIX_ROW_PADDING != 0) { + total_size += ggml_row_size(slice_type, MATRIX_ROW_PADDING - ne0 % MATRIX_ROW_PADDING); + } + } + return total_size; +} + +static bool ggml_backend_cuda_split_buffer_type_is_host(ggml_backend_buffer_type_t buft) { + return false; + GGML_UNUSED(buft); +} + +static ggml_backend_buffer_type_i ggml_backend_cuda_split_buffer_type_interface = { + /* .get_name = */ ggml_backend_cuda_split_buffer_type_get_name, + /* .alloc_buffer = */ ggml_backend_cuda_split_buffer_type_alloc_buffer, + /* .get_alignment = */ ggml_backend_cuda_split_buffer_type_get_alignment, + /* .get_max_size = */ NULL, + /* .get_alloc_size = */ ggml_backend_cuda_split_buffer_type_get_alloc_size, + /* .is_host = */ ggml_backend_cuda_split_buffer_type_is_host, +}; + +ggml_backend_buffer_type_t ggml_backend_cuda_split_buffer_type(int main_device, const float * tensor_split) { + GGML_UNUSED(main_device); + static std::mutex mutex; + std::lock_guard lock(mutex); + static std::map, struct ggml_backend_buffer_type> buft_map; + std::array tensor_split_arr = {}; + bool all_zero = tensor_split == nullptr || std::all_of(tensor_split, tensor_split + GGML_CUDA_MAX_DEVICES, [](float x){ return x==0.0f; }); + if (all_zero) { + // use default split from ggml_cuda_info (balanced) + // fall back to equal split if default not set + for (int i=0;i0) { + // try to use real default if available + for (int i=0;isecond; + ggml_backend_cuda_split_buffer_type_context * ctx = new ggml_backend_cuda_split_buffer_type_context(); + ctx->main_device = main_device; + ctx->tensor_split = tensor_split_arr; + ctx->name = GGML_CUDA_NAME "_Split"; + struct ggml_backend_buffer_type buft = { + /* .iface = */ ggml_backend_cuda_split_buffer_type_interface, + /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cuda_reg(), main_device), + /* .context = */ ctx, + }; + auto result = buft_map.emplace(tensor_split_arr, buft); + return &result.first->second; +} + // Communication context for multi-GPU AllReduce during tensor parallelism. // // Created once per meta backend instance. Resources for the selected mode @@ -2150,6 +2500,47 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * } static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct ggml_tensor * dst) { + // PR104 mixed-quant: generic gather for any split src (SCALE, etc. also need it) + // Disable CUDA graph for split src - host staging and peer copies are not graph-capture safe + { + cudaStreamCaptureStatus cap_status; + if (cudaStreamIsCapturing(ctx.stream(), &cap_status) == cudaSuccess && cap_status == cudaStreamCaptureStatusActive) { + for (int i = 0; i < 4; ++i) { + ggml_tensor * src = dst->src[i]; + if (src && src->buffer && ggml_backend_buffer_is_cuda_split(src->buffer)) { + return false; + } + } + } + } + ggml_tensor * orig_src[4] = {dst->src[0], dst->src[1], dst->src[2], dst->src[3]}; + ggml_tensor src_view[4]; + bool is_split[4] = {}; + for (int i = 0; i < 4; ++i) { + ggml_tensor * src = orig_src[i]; + if (src && src->buffer && ggml_backend_buffer_is_cuda_split(src->buffer)) { + is_split[i] = true; + ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *) src->extra; + GGML_LOG_INFO("PR104 split gather: op=%s src%d type=%s contiguous=%p dev=%d ctx=%d\n", + ggml_op_desc(dst), i, ggml_type_name(src->type), extra->contiguous_data, extra->contiguous_device, ctx.device); + // use cached contiguous copy (filled at load time, on main_device) + // for now we require contiguous on same device as ctx (main_device is 0) + // if mismatch, we would need peer copy, but for LLM main is 0 and ctx is 0 + if (!extra->contiguous_data) { + GGML_LOG_ERROR("PR104 split gather: contiguous is null for %s\n", ggml_op_desc(dst)); + GGML_ABORT("contiguous null"); + } + if (extra->contiguous_device != ctx.device) { + GGML_LOG_WARN("PR104 split gather: contiguous on %d but ctx %d, peer copy needed\n", extra->contiguous_device, ctx.device); + // for now, just use it anyway (peer access may be needed, but try) + } + src_view[i] = *src; + src_view[i].data = extra->contiguous_data; + // keep original buffer/extra but data now points to contiguous + dst->src[i] = &src_view[i]; + } + } + switch (dst->op) { case GGML_OP_ARGMAX: ggml_cuda_argmax(ctx, dst); @@ -2498,9 +2889,14 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg ggml_cuda_lightning_indexer(ctx, dst); break; default: + // restore original src pointers if we had split gather + for (int i = 0; i < 4; ++i) if (is_split[i]) dst->src[i] = orig_src[i]; return false; } + // restore original src pointers after gather + for (int i = 0; i < 4; ++i) if (is_split[i]) dst->src[i] = orig_src[i]; + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { GGML_LOG_ERROR("%s: %s failed\n", __func__, ggml_op_desc(dst)); @@ -5608,7 +6004,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g static bool ggml_backend_cuda_device_supports_buft(ggml_backend_dev_t dev, ggml_backend_buffer_type_t buft) { ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *) dev->context; const bool integrated = ggml_cuda_info().devices[dev_ctx->device].integrated; - return (ggml_backend_buft_is_cuda(buft) && buft->device == dev) || (integrated && ggml_backend_buft_is_cuda_host(buft)); + return ((ggml_backend_buft_is_cuda(buft) || ggml_backend_buft_is_cuda_split(buft)) && buft->device == dev) || (integrated && ggml_backend_buft_is_cuda_host(buft)); } static int64_t get_op_batch_size(const ggml_tensor * op) { @@ -5780,6 +6176,9 @@ static void * ggml_backend_cuda_reg_get_proc_address(ggml_backend_reg_t reg, con if (strcmp(name, "ggml_backend_get_features") == 0) { return (void *)ggml_backend_cuda_get_features; } + if (strcmp(name, "ggml_backend_split_buffer_type") == 0) { + return (void *)ggml_backend_cuda_split_buffer_type; + } return nullptr; } From 40ea56978cc4989d3f01c010c7082da491b187df Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 21:14:54 +0700 Subject: [PATCH 5/8] cuda: harden split-buffer compute path - gate CUDA graph capture off for graph nodes with split-buffer srcs (mid-capture false returns were silently skipped in release builds, leaving malformed captured graphs) - harden the contiguous-gather guard: abort on contiguous copy on the wrong compute device instead of reading across devices without peer access - gate fusion windows in try_fuse so fused kernels (which bypass the per-op gather and read split tensor->data directly) never dispatch with split-backed srcs - document why split get_base must stay a pseudo pointer --- ggml/src/ggml-cuda/ggml-cuda.cu | 87 +++++++++++++++++++++++++++------ 1 file changed, 71 insertions(+), 16 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 8b5c0b386a2a..59ba1e3be84f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1118,6 +1118,10 @@ static void ggml_backend_cuda_split_buffer_free_buffer(ggml_backend_buffer_t buf } static void * ggml_backend_cuda_split_buffer_get_base(ggml_backend_buffer_t buffer) { + // Split storage never aliases buffer base memory: per-device slices + the cached contiguous copy + // live in tensor extra. The fake base only exists to satisfy the assertions in ggml_tallocr_alloc / + // ggml_backend_tensor_alloc (addr range >= base and <= base + size); no code must dereference + // split tensor->data directly - compute paths must use the cached contiguous copy instead. return (void *)0x1000; GGML_UNUSED(buffer); } @@ -2501,14 +2505,16 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct ggml_tensor * dst) { // PR104 mixed-quant: generic gather for any split src (SCALE, etc. also need it) - // Disable CUDA graph for split src - host staging and peer copies are not graph-capture safe + // Split-src nodes must never be captured: the gather overrides src data pointers with + // evaluation-local views, so a replayed graph would read dead pointers. graph_check_compability + // disables capture for split-src nodes; this guard is a backstop. { cudaStreamCaptureStatus cap_status; if (cudaStreamIsCapturing(ctx.stream(), &cap_status) == cudaSuccess && cap_status == cudaStreamCaptureStatusActive) { for (int i = 0; i < 4; ++i) { ggml_tensor * src = dst->src[i]; if (src && src->buffer && ggml_backend_buffer_is_cuda_split(src->buffer)) { - return false; + GGML_ABORT("split-buffer src during CUDA graph capture - graph compat gate missed this node"); } } } @@ -2531,8 +2537,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg GGML_ABORT("contiguous null"); } if (extra->contiguous_device != ctx.device) { - GGML_LOG_WARN("PR104 split gather: contiguous on %d but ctx %d, peer copy needed\n", extra->contiguous_device, ctx.device); - // for now, just use it anyway (peer access may be needed, but try) + GGML_LOG_ERROR("PR104 split gather: contiguous on %d but ctx %d, cross-device read without peer access\n", + extra->contiguous_device, ctx.device); + GGML_ABORT("split contiguous copy on wrong device"); } src_view[i] = *src; src_view[i].data = extra->contiguous_data; @@ -3061,6 +3068,18 @@ static bool ggml_cuda_graph_check_compability(ggml_cgraph * cgraph) { } } + // PR104 mixed-quant: split-src ops must not be captured - per-node gather replacements in + // compute_forward decay after this evaluation, so replayed graphs would read stale pointers + for (int j = 0; j < GGML_MAX_SRC; j++) { + const ggml_tensor * src = node->src[j]; + if (src && src->buffer && ggml_backend_buffer_is_cuda_split(src->buffer)) { + use_cuda_graph = false; +#ifndef NDEBUG + GGML_LOG_DEBUG("%s: disabling CUDA graphs due to split buffer src on node %s\n", __func__, node->name); +#endif + } + } + if (!use_cuda_graph) { break; } @@ -3667,6 +3686,22 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, GGML_ASSERT(unary_ops.size() == num_unary); #endif + // PR104 mixed-quant: fused kernels read tensor->data of the participating nodes' srcs directly + // and bypass the per-op gather in compute_forward that substitutes the cached contiguous copy for + // split tensors. Refuse any fusion window covered by a split-backed src. + for (int k = node_idx; k < node_idx + (int) ops.size() && k < cgraph->n_nodes; ++k) { + const ggml_tensor * n = cgraph->nodes[k]; + if (!n || (n->buffer && ggml_backend_buffer_is_cuda_split(n->buffer))) { + return false; + } + for (int j = 0; j < GGML_MAX_SRC; ++j) { + const ggml_tensor * s = n->src[j]; + if (s && s->buffer && ggml_backend_buffer_is_cuda_split(s->buffer)) { + return false; + } + } + } + const auto is_equal = [](const std::initializer_list & list1, const std::initializer_list & list2) { return std::equal(list1.begin(), list1.end(), list2.begin(), list2.end()); @@ -3917,13 +3952,29 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph return 0; } + // PR104 mixed-quant: fused kernels read tensor->data of the participating nodes' srcs directly + // and bypass the per-op gather in compute_forward that substitutes the cached contiguous copy for + // split tensors. Returns true when the nodes in [node_idx, node_idx + span) are free of split refs. + auto span_open = [cgraph](int node_idx, int span) -> bool { + for (int k = node_idx; k < node_idx + span && k < cgraph->n_nodes; ++k) { + ggml_tensor * n = cgraph->nodes[k]; + if (!n) return false; + if (n->buffer && ggml_backend_buffer_is_cuda_split(n->buffer)) return false; + for (int j = 0; j < GGML_MAX_SRC; ++j) { + const ggml_tensor * s = n->src[j]; + if (s && s->buffer && ggml_backend_buffer_is_cuda_split(s->buffer)) return false; + } + } + return true; + }; + ggml_tensor * node = cgraph->nodes[i]; if (node->op == GGML_OP_MUL) { ggml_cuda_moe_weighted_reduction_match match; if (ggml_cuda_match_moe_weighted_reduction(cgraph, i, match)) { const int output_idx = i + match.node_count - 1; - if (ggml_cuda_check_fusion_memory_ranges(cgraph, i, match.node_count, &output_idx, 1)) { + if (ggml_cuda_check_fusion_memory_ranges(cgraph, i, match.node_count, &output_idx, 1) && span_open(i, match.node_count)) { ggml_cuda_op_moe_weighted_reduction( *cuda_ctx, match.experts, match.expert_scale, match.weights, match.dst); return match.node_count - 1; @@ -3935,7 +3986,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph if (node->op == GGML_OP_GATED_DELTA_NET) { ggml_cuda_gated_delta_net_fused_cache fused_state_cpy; const int nodes_to_skip = ggml_cuda_try_gdn_cache_fusion(cgraph, i, fused_state_cpy); - if (nodes_to_skip > 0) { + if (nodes_to_skip > 0 && span_open(i, nodes_to_skip + 1)) { #ifdef GGML_CUDA_DEBUG GGML_LOG_INFO("%s: fused gated_delta_net snapshot copies for %s (skipped %d nodes)\n", __func__, node->name, nodes_to_skip); @@ -3998,7 +4049,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && ggml_cuda_should_use_topk_moe(node, logits, weights, ids) && - ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { + ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true) && + span_open(i, (int) ops.size())) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4013,7 +4065,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph int out_nodes[2] = { i + 1, i + 5 }; if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && ggml_cuda_should_use_topk_moe(softmax, logits, weights, ids) && - ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { + ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true) && + span_open(i, (int) ops.size())) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4071,7 +4124,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph const bool contig_ok = ggml_is_contiguous(x) && ggml_is_contiguous(add) && ggml_is_contiguous(a) && ggml_is_contiguous(inv_b); - if (types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x) { + if (types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x && span_open(i, 5)) { ggml_cuda_op_snake_fused(*cuda_ctx, x, a, inv_b, add); return 4; } @@ -4097,7 +4150,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph n_fuse++; - if (n_fuse > 1) { + if (n_fuse > 1 && span_open(i, n_fuse + 1)) { ggml_tensor fused_node; memcpy(&fused_node, node, sizeof(ggml_tensor)); for (int j = 0; j < n_fuse - 1; ++j) { @@ -4202,7 +4255,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph const int n_ops = with_bias ? 7 : 5; if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || - !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) { + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1) || + !span_open(i, n_ops)) { continue; } @@ -4498,10 +4552,11 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } - if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || - !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) { - continue; - } + if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1) || + !span_open(i, n_ops)) { + continue; + } ggml_tensor * mm_node = cgraph->nodes[i]; ggml_tensor * scale_node = op == GGML_OP_MUL_MAT ? cgraph->nodes[i + 1] : cgraph->nodes[i + 4]; @@ -4556,7 +4611,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; - if (!ggml_can_fuse(cgraph, i, { op, bias_op })) { + if (!ggml_can_fuse(cgraph, i, { op, bias_op }) || !span_open(i, 2)) { continue; } From 23e204be4233e14e935ba6041a21349e3b9a70d6 Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Thu, 10 Sep 2026 23:08:57 +0700 Subject: [PATCH 6/8] docs: PR104.2 compute crash root-caused and fixed, 9B validated; 27B needs per-op split dispatch --- docs/arms/pr104-mixed-quant-row-sharding.md | 28 +++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md index 6c6cb41087b6..74e588b3eba6 100644 --- a/docs/arms/pr104-mixed-quant-row-sharding.md +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -333,3 +333,31 @@ above), `nvidia-smi` before/after 1/15712→15847 MiB. Commit on **Rig handoff:** `nvidia-smi` `1 MiB / 16311` `1 MiB / 12288` `0%`, `pod_llama-baseline Exited` (not `Running`), no active `llama-server`/`rpc` (only `1q3ry0vb` defunct zombies from prior runs, `1 MiB` free). **Not restarting production** — ae62ab1e needs rig for PR105.0. Standing down. **Commit:** `ggml/src/ggml-cuda/common.cuh` (`slice_type` + `contiguous_data`), `ggml/src/ggml-cuda/ggml-cuda.cu` (split buffer + `contiguous` cache + `compute_forward` graph guard), docs appended. To be pushed `ddvnguyen fork/pr104-mixed-quant-arm` (updates PR #113). + +## PR104.2 Progress 2 — compute crash root-caused and fixed; 27B needs per-op split dispatch — 2026-09-10 (session 2) + +**Pickup state:** storage half already validated (session 1). This session root-caused and fixed the compute-path crash, validated 9B end-to-end decode on the split path, and mapped the remaining 27B work. + +**Root cause (confirmed by code-path analysis + compute-sanitizer + empirical A/B):** three interacting defects in the split-buffer compute wiring, all in `ggml/src/ggml-cuda/ggml-cuda.cu`: + +1. **CUDA-graph capture corruption.** `compute_forward` returned `false` mid-capture to disable capture for split srcs — but `GGML_ASSERT(ok)` in `ggml_cuda_graph_evaluate_and_capture` is a no-op in Release builds, so the node was silently skipped and the stream capture still continued/completed, producing malformed captured graphs. With graph capture active the `Illegal memory access` surfaced at the next kernel launch (`scale` right after the first split MUL_MAT — the reported signature). Fix: gate capture off up-front in `ggml_cuda_graph_check_compability` (any node with a split-buffer src → `use_cuda_graph=false`, same pattern as the `MUL_MAT_ID` tag), and turned the mid-capture guard into a hard `GGML_ABORT` backstop. +2. **Silent cross-device read in the gather.** `contiguous_data` is allocated on `buft_ctx->main_device`; if the compute ctx device differs the old code warned ("uses it anyway") and read a foreign-device pointer without peer access → illegal access. Now: hard `GGML_ABORT` on `contiguous_device != ctx.device` (evidence shows dev=0/ctx=0 always match on this rig; abort did not fire in validation). +3. **Fused-op bypass of the gather.** All MUL_MAT-family fusions (`ggml_cuda_try_fuse`: fused MMVQ+GLU, mul_mat+add, 5/7/11/13-op MUL_MAT+REPEAT subgraphs, fused add/mul chains, rms_norm+mul+rope, snake, topk-moe, moe-weighted-reduction, gdn cache) dispatch fused kernels directly with `cgraph->nodes[]` srcs and read `tensor->data` — which for split tensors is the fake base `0x1000+offset` (`get_base`) — bypassing the per-op gather entirely. Any fusion window covering a split-backed src is now refused: a split-ref check inside `ggml_cuda_can_fuse` (span = its own op pattern) plus per-dispatch `span_open(i, span)` gates in `try_fuse` for the `ggml_can_fuse_subgraph`-based families (snake, topk-moe both variants, moe-weighted-reduction, gdn-cache, gate/glu/up 5/7/11/13-op subgraphs, mul_mat+add). + +`get_base` stays a pseudo pointer (documented in place: tallocr range asserts require a non-null base; split tensors must never dereference `tensor->data`). + +**Validation (9B probe, `Qwen3.5-9B-Q4_K_M`, `-sm row -ts 38,27`, `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k`):** + +- `compute-sanitizer memcheck` pre-fix surfaced the live signature `CUDA Stream does not belong to the expected context` at `ggml_backend_cuda_synchronize` during the fit-probe context teardown — consistent with the capture/skip corruption class above (no attempt to fix that specific fit-probe teardown; the production crash itself is covered by fix 1/3). +- Post-fix boot: health 200, VRAM 8089/3109 MiB (slices + contiguous copies + KV), `PR104 split gather` fired **852×** during warmup decode with `dev=0 ctx=0` (device match; activation shows `MUL_MAT` split src0 `q5_K` and `MUL` split src1 f32 norm weights), completion `"The capital of France is Paris"` — **correct tokens**, `,graphs reused`=..., decode **41.7 t/s**, pp **155 t/s** (no MTP). The device-mismatch abort never fired. Capture-disable path logged `disabling CUDA graphs` **4685×** without tripping the backstop abort. +- Vanilla row-split (no env) also works pod-stop-checked (earlier 25s boot observed 8.4/2.4 GiB mid-load before shell-timeout kill; load-time requant confirmed active). + +**27B blocker (new, empirical):** 27B mixed-quant boot dies at load with `CUDA error: out of memory` in split `init_tensor`. Cause: the per-tensor contiguous copies double the weights on the main device (~5.3 GiB extra for 9B observed; for 27B that is an 18.4 GiB copy on a 16 GB card). NOTE: session-1's doc claims of 27B row-split VRAM 8283/9001 predate/contradict this test — treat those numbers as unreproduced. + +**Remaining work (next session, well-defined):** +1. Port SYCL/old-CUDA per-op row-split `ggml_cuda_op_mul_mat` (the complete mechanism lives at git history `ae8de6d50^:ggml/src/ggml-cuda.cu:1345` — `MUL_MAT_SRC1_COL_STRIDE=128`, `ggml_cuda_Memcpy2DPeerAsync` via `cudaMemcpy3DPeerAsync`, `get_mmq_x_max_host` J-padding staging, per-device events/stream loops, dst partial `Memcpy2DPeerAsync` stitch). Modern counterpart for MMVQ already exists (`ggml_cuda_op_mul_mat_vec_q` in `mmvq.cu` with row_low/high/dst_dd_i signature); needs the MMQ per-op shim in `mmq.cu` via `mmq_args`+`ggml_cuda_mmq_get_J_max`, and MMVF/CUBLAS fallback policy. +2. `compute_forward`: skip the gather for MUL_MAT (dispatch handles split directly); keep the gather for small split tensors consumed by other ops. +3. `init_tensor`/`set_tensor`: contiguous copy only under a bytes threshold (e.g. `ggml_nbytes(tensor) <= 256 MiB`) so MUL_MAT weight tensors stay slices-only. +4. Keep: capture gate, fusion gates, gather for non-MUL_MAT split sources. Then 27B mixed-quant boot + `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k` MTP production run, `test-suite.sh `, decode t/s vs ≥45 bar, PR104.0 projection comparison, fork push. + +**Commit:** fix 1-3 + validation is `40ea56978 cuda: harden split-buffer compute path` on `fork/pr104-mixed-quant-arm` (PR #113). Production pod lifecycle protocol followed this session: `podman pod stop pod_llama-baseline` before test boots (VRAM drained 15847/11911 → 1/1 MiB), `podman pod start` restored — health 200 on :18081, VRAM back to 15847/11911 MiB. From a38976d983f1514f1ca447fa9826beb1cafc1479 Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Fri, 11 Sep 2026 00:48:32 +0700 Subject: [PATCH 7/8] cuda: per-op row-split MUL_MAT dispatch for mixed-quant split buffers --- ggml/src/ggml-cuda/ggml-cuda.cu | 258 ++++++++++++++++++++++++++++++-- ggml/src/ggml-cuda/mmq.cu | 57 +++++++ ggml/src/ggml-cuda/mmq.cuh | 7 + ggml/src/ggml-cuda/mmvq.cu | 9 +- ggml/src/ggml-cuda/mmvq.cuh | 2 +- 5 files changed, 314 insertions(+), 19 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 59ba1e3be84f..0208058e9007 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1164,22 +1164,36 @@ static enum ggml_status ggml_backend_cuda_split_buffer_init_tensor(ggml_backend_ } } // for devices without rows, keep slice_type as COUNT (fallback to tensor->type in dispatch) - // also alloc contiguous copy on main_device for compute (avoids per-op gather, graph-safe) + // also alloc contiguous copy on main_device for compute of ops without a native split path + // (PR104.2: gated by size so weight tensors consumed by the split MUL_MAT dispatch stay slices-only) { int main_id = buft_ctx->main_device; if (main_id < 0 || main_id >= ggml_backend_cuda_get_device_count()) main_id = 0; - size_t cont_size = ggml_nbytes(tensor); - if (ne0 % MATRIX_ROW_PADDING != 0) { - cont_size += ggml_row_size(tensor->type, MATRIX_ROW_PADDING - ne0 % MATRIX_ROW_PADDING); - } - ggml_cuda_set_device(main_id); - void * cbuf = nullptr; - CUDA_CHECK(ggml_cuda_device_malloc(&cbuf, cont_size, main_id)); - if (cont_size > ggml_nbytes(tensor)) { - CUDA_CHECK(cudaMemset((char*)cbuf + ggml_nbytes(tensor), 0, cont_size - ggml_nbytes(tensor))); + static size_t contig_max_bytes = [] { + const char * env = getenv("GGML_SPLIT_CONTIG_MAX_BYTES"); + // PR104.2: contiguous copies are only needed by ops without a native split path + // (norm/bias-scale MUL etc.); weight tensors stay slices-only for the split MUL_MAT dispatch. + // 16 MiB keeps all small f32 norm/rope weights and drops mid-size layer weights. + return env ? (size_t) std::atoll(env) : (size_t) 16*1024*1024; + }(); + if (ggml_nbytes(tensor) <= contig_max_bytes) { + size_t cont_size = ggml_nbytes(tensor); + if (ne0 % MATRIX_ROW_PADDING != 0) { + cont_size += ggml_row_size(tensor->type, MATRIX_ROW_PADDING - ne0 % MATRIX_ROW_PADDING); + } + fprintf(stderr, "PR104 split contig: %s type=%s size=%zu cont=%zu\n", tensor->name, ggml_type_name(tensor->type), (size_t) ggml_nbytes(tensor), cont_size); + ggml_cuda_set_device(main_id); + void * cbuf = nullptr; + CUDA_CHECK(ggml_cuda_device_malloc(&cbuf, cont_size, main_id)); + if (cont_size > ggml_nbytes(tensor)) { + CUDA_CHECK(cudaMemset((char*)cbuf + ggml_nbytes(tensor), 0, cont_size - ggml_nbytes(tensor))); + } + extra->contiguous_data = cbuf; + extra->contiguous_device = main_id; + } else { + extra->contiguous_data = nullptr; + extra->contiguous_device = main_id; } - extra->contiguous_data = cbuf; - extra->contiguous_device = main_id; } tensor->extra = extra; return GGML_STATUS_SUCCESS; @@ -1231,13 +1245,15 @@ static void ggml_backend_cuda_split_buffer_set_tensor(ggml_backend_buffer_t buff // padding already zeroed in init GGML_UNUSED(ne0); } - // also copy full tensor to contiguous on main_device (original type) + // also copy full tensor to contiguous on main_device (original type), when the tensor qualifies { int main_id = extra->contiguous_device; if (main_id < 0) main_id = buft_ctx->main_device; if (main_id < 0 || main_id >= ggml_backend_cuda_get_device_count()) main_id = 0; - ggml_cuda_set_device(main_id); - CUDA_CHECK(cudaMemcpy(extra->contiguous_data, data, ggml_nbytes(tensor), cudaMemcpyHostToDevice)); + if (extra->contiguous_data) { + ggml_cuda_set_device(main_id); + CUDA_CHECK(cudaMemcpy(extra->contiguous_data, data, ggml_nbytes(tensor), cudaMemcpyHostToDevice)); + } } GGML_UNUSED(buffer); } @@ -2259,9 +2275,215 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) { return use_mul_mat_vec_q; } +// PR104.2: row-split MUL_MAT. Each device runs a native quantized matmul over its own weights slice, so no +// per-op contiguous gather is needed and devices compute concurrently. Port of the removed per-op +// ggml_cuda_op_mul_mat (SYCL ggml_sycl_op_mul_mat), with per-device chunked src1 staging. +#define PR104_MUL_MAT_SRC1_COL_STRIDE 128 + +// cudaMemcpy2DAsync may fail with copies between vmm pools of different devices, so use cudaMemcpy3DPeerAsync +static cudaError_t ggml_cuda_Memcpy2DPeerAsync( + void * dst, int dstDevice, size_t dpitch, void * src, int srcDevice, size_t spitch, size_t width, size_t height, cudaStream_t stream) { + + cudaMemcpy3DPeerParms p = {}; + p.dstDevice = dstDevice; + p.dstPtr = make_cudaPitchedPtr(dst, dpitch, dpitch, height); + p.srcDevice = srcDevice; + p.srcPtr = make_cudaPitchedPtr(src, spitch, spitch, height); + p.extent = make_cudaExtent(width, height, 1); + return cudaMemcpy3DPeerAsync(&p, stream); +} + +static void ggml_cuda_mul_mat_split(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(ggml_backend_buffer_is_cuda_split(src0->buffer)); + + GGML_TENSOR_BINARY_OP_LOCALS; + + ggml_backend_cuda_split_buffer_type_context * buft_ctx = (ggml_backend_cuda_split_buffer_type_context *) src0->buffer->buft->context; + ggml_tensor_extra_gpu * src0_extra = (ggml_tensor_extra_gpu *) src0->extra; + GGML_ASSERT(src0_extra); + + GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(!dst->src[2]); // no mul_mat ids + GGML_ASSERT(nb00 == ggml_type_size(src0->type)); + GGML_ASSERT(nb10 == ggml_type_size(src1->type)); + GGML_ASSERT(nb0 == sizeof(float)); + + // split weights are 2D; broadcast (ne12 != ne02) src1 and split src1/dst are not supported + GGML_ASSERT(ne02 == 1 && ne03 == 1); + GGML_ASSERT(ne12 == 1 && ne13 == 1); + GGML_ASSERT(!ggml_backend_buffer_is_cuda_split(src1->buffer)); + GGML_ASSERT(!ggml_backend_buffer_is_cuda_split(dst->buffer)); + GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous(dst)); + GGML_ASSERT(dst->nb[1] == ne0*sizeof(float)); // dst stitch copies full-size rows + + const int cc = ggml_cuda_info().devices[ctx.device].cc; + const ggml_type type_src0 = src0->type; + GGML_ASSERT(ggml_is_quantized(type_src0)); + + const bool use_mmvq = ggml_cuda_should_use_mmvq(type_src0, cc, ne11); + const bool use_mmq = !use_mmvq && ggml_cuda_should_use_mmq(type_src0, cc, ne11, /*n_experts=*/0); + if (!use_mmvq && !use_mmq) { + GGML_ABORT("PR104: split-buffer MUL_MAT needs a quantized src0 (mmvq/mmq), got %s", ggml_type_name(type_src0)); + } + + const int dev_src1 = ((ggml_backend_cuda_buffer_context *) src1->buffer->context)->device; + + struct dev_data { + int64_t row_low = 0; + int64_t row_high = 0; + + ggml_cuda_pool_alloc src1_ddf_alloc; + ggml_cuda_pool_alloc dst_dd_alloc; + + float * dst_dd = nullptr; + float * src1_ddf = nullptr; + }; + + dev_data dev[GGML_CUDA_MAX_DEVICES]; + + int used_devices = 0; + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + dev[id].row_low = 0; + dev[id].row_high = ne01; + + // for multi GPU, get the row boundaries from the tensor split and round to mul_mat_q tile sizes + get_row_split_cuda(&dev[id].row_low, &dev[id].row_high, src0, buft_ctx->tensor_split, id); + + if (dev[id].row_low == dev[id].row_high) { + continue; + } + + used_devices++; + + // the main device writes its result directly into dst (per-op kernels offset by row_low), + // every other device needs a token-major dst staging buffer + const bool dst_on_device = id == ctx.device; + const int64_t row_diff = dev[id].row_high - dev[id].row_low; + + ggml_cuda_set_device(id); + cudaStream_t stream = ctx.stream(id, 0); + + dev[id].src1_ddf = dev[id].src1_ddf_alloc.alloc(ctx.pool(id), ne11*ne10); + + if (dst_on_device) { + dev[id].dst_dd = (float *) dst->data; + } else { + dev[id].dst_dd = dev[id].dst_dd_alloc.alloc(ctx.pool(id), row_diff*ne1); + } + } + GGML_ASSERT(used_devices > 0); + + // if multiple devices are used, they need to wait for the main device's input data to be ready + if (used_devices > 1) { + ggml_cuda_set_device(ctx.device); + CUDA_CHECK(cudaEventRecord(src0_extra->events[ctx.device][0], ctx.stream())); + } + + const int64_t src1_col_stride = used_devices > 1 ? PR104_MUL_MAT_SRC1_COL_STRIDE : ne11; + for (int64_t src1_col_0 = 0; src1_col_0 < ne11; src1_col_0 += src1_col_stride) { + const int64_t is = used_devices > 1 ? (src1_col_0/src1_col_stride) % GGML_CUDA_MAX_STREAMS : 0; + const int64_t src1_ncols = src1_col_0 + src1_col_stride > ne11 ? ne11 - src1_col_0 : src1_col_stride; + const int64_t ne10_padded = GGML_PAD(ne10, MATRIX_ROW_PADDING); + + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + if (dev[id].row_low == dev[id].row_high) { + continue; + } + + const bool dst_on_device = id == ctx.device; + const int64_t row_diff = dev[id].row_high - dev[id].row_low; + + ggml_cuda_set_device(id); + cudaStream_t stream = ctx.stream(id, is); + + // wait for main GPU data if necessary + if (used_devices > 1 && (id != ctx.device || is != 0)) { + CUDA_CHECK(cudaStreamWaitEvent(stream, src0_extra->events[ctx.device][0], 0)); + } + + // for split tensors the src0 slice is already on the computing device + char * src0_dd_i = (char *) src0_extra->data_device[id]; + + // the main device can write its result directly into dst (the per-op kernels offset the + // written rows by row_low), all other devices write into a staging buffer + float * dst_dd_i = dev[id].dst_dd + (dst_on_device ? src1_col_0*ne0 + dev[id].row_low : src1_col_0*row_diff); + + // stage the src1 columns on a device that does not own src1 + float * src1_ddf_i = nullptr; + if (id != dev_src1) { + src1_ddf_i = dev[id].src1_ddf + src1_col_0*ne10; + CUDA_CHECK(cudaMemcpyPeerAsync(src1_ddf_i, id, + ((const float *) src1->data) + src1_col_0*ne10, dev_src1, + src1_ncols*ne10*sizeof(float), stream)); + } else { + src1_ddf_i = (float *) src1->data + src1_col_0*ne10; + } + + const bool use_mmvq = ggml_cuda_should_use_mmvq(type_src0, cc, ne11); + // PR104.2: the slice on this device may be re-quantized to a different type + const ggml_type type_slice = src0_extra->slice_type[id] != GGML_TYPE_COUNT ? src0_extra->slice_type[id] : type_src0; + + if (use_mmvq) { + ggml_cuda_pool_alloc src1_ddq_alloc; + char * src1_ddq_i = src1_ddq_alloc.alloc(ctx.pool(id), src1_ncols*ne10_padded*sizeof(block_q8_1)/QK8_1); + quantize_row_q8_1_cuda(src1_ddf_i, nullptr, src1_ddq_i, type_slice, ne10, ne10, 1, 1, + ne10_padded, src1_ncols, 1, 1, stream); + CUDA_CHECK(cudaGetLastError()); + + ggml_cuda_op_mul_mat_vec_q(ctx, src0, src1, dst, src0_dd_i, src1_ddf_i, src1_ddq_i, + dst_dd_i, dev[id].row_low, dev[id].row_high, src1_ncols, ne10_padded, stream, type_slice, ctx.pool(id)); + } else { + ggml_cuda_op_mul_mat_q(ctx, src0, src1, dst, src0_dd_i, src1_ddf_i, dst_dd_i, + dev[id].row_low, dev[id].row_high, src1_ncols, ne10_padded, stream, type_slice, ctx.pool(id)); + } + CUDA_CHECK(cudaGetLastError()); + + // copy dst to the main device if necessary + if (!dst_on_device) { + // src0 is stored row-major, so the outputs of matrix multiplications cannot simply be + // concatenated for >1 GPU: they need to be copied to the correct slice in ne0 = dst rows + float * dhf_dst_i = (float *) dst->data + src1_col_0*ne0 + dev[id].row_low; + CUDA_CHECK(ggml_cuda_Memcpy2DPeerAsync( + dhf_dst_i, ctx.device, ne0*sizeof(float), dst_dd_i, id, + row_diff*sizeof(float), row_diff*sizeof(float), src1_ncols, stream)); + } + + // add event for the main device to wait on until other device is done + if (used_devices > 1 && (id != ctx.device || is != 0)) { + CUDA_CHECK(cudaEventRecord(src0_extra->events[id][is], stream)); + } + } + } + + // main device waits for all other devices to be finished + if (used_devices > 1) { + int64_t is_max = (ne11 + PR104_MUL_MAT_SRC1_COL_STRIDE - 1) / PR104_MUL_MAT_SRC1_COL_STRIDE; + is_max = std::min(is_max, (int64_t) GGML_CUDA_MAX_STREAMS); + + ggml_cuda_set_device(ctx.device); + for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { + if (dev[id].row_low == dev[id].row_high) { + continue; + } + for (int64_t isl = 0; isl < is_max; ++isl) { + CUDA_CHECK(cudaStreamWaitEvent(ctx.stream(), src0_extra->events[id][isl], 0)); + } + } + } +} + static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_TENSOR_BINARY_OP_LOCALS + // PR104.2: split-buffer src0 is dispatched as a partially-mapped native matmul over the per-device + // weights slices, no gather needed + if (src0->buffer && ggml_backend_buffer_is_cuda_split(src0->buffer)) { + ggml_cuda_mul_mat_split(ctx, src0, src1, dst); + return; + } + const int32_t hint = ggml_get_op_params_i32(dst, 1); if (hint == GGML_HINT_SRC0_IS_HADAMARD && ggml_cuda_op_fwht(ctx, src1, dst)) { return; @@ -2522,9 +2744,15 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg ggml_tensor * orig_src[4] = {dst->src[0], dst->src[1], dst->src[2], dst->src[3]}; ggml_tensor src_view[4]; bool is_split[4] = {}; + // PR104.2: for MUL_MAT the split src0 dispatch handles the split directly, do not gather it + const bool mul_mat_split_src0 = dst->op == GGML_OP_MUL_MAT && orig_src[0] && orig_src[0]->buffer && + ggml_backend_buffer_is_cuda_split(orig_src[0]->buffer); for (int i = 0; i < 4; ++i) { ggml_tensor * src = orig_src[i]; if (src && src->buffer && ggml_backend_buffer_is_cuda_split(src->buffer)) { + if (mul_mat_split_src0 && i == 0) { + continue; + } is_split[i] = true; ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *) src->extra; GGML_LOG_INFO("PR104 split gather: op=%s src%d type=%s contiguous=%p dev=%d ctx=%d\n", diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 7fb4401489c9..0f1478211dde 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -256,6 +256,63 @@ void ggml_cuda_mul_mat_q( ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } +// PR104.2: per-op MUL_MAT for CUDA row-split buffers (port of removed ggml_cuda_op_mul_mat_mmq / +// SYCL ggml_sycl_op_mul_mat). +// src0 is split: src0_dd_i is this device's quantized slice [row_low, row_high). +// src1_ddf_i is a tightly packed f32 chunk of src1_ncols rows (ne10 elements each), already staged on the +// current device. dst_dd_i is either a direct view into the real dst (main device, global rows) or a +// token-major partial (pitch row_diff). No source quantization needed: each call quantizes the chunk. +void ggml_cuda_op_mul_mat_q( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, + const char * src0_dd_i, const float * src1_ddf_i, float * dst_dd_i, + const int64_t row_low, const int64_t row_high, const int64_t src1_ncols, + const int64_t src1_padded_row_size, cudaStream_t stream, const ggml_type type_slice, ggml_cuda_pool & pool_slice) { + + const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + const ggml_type type_x = type_slice == GGML_TYPE_COUNT ? src0->type : type_slice; + + const int64_t ne00 = src0->ne[0]; + const int64_t row_diff = row_high - row_low; + + const bool fallback = src0->ne[1] % 128 != 0; + const bool use_native_fp4 = blackwell_mma_available(cc) && (type_x == GGML_TYPE_MXFP4 || type_x == GGML_TYPE_NVFP4); + if (use_native_fp4) { + GGML_ABORT("PR104.2: fp4 mmq split not implemented"); + } + + const int64_t ne10_padded = GGML_PAD(src1->ne[0], MATRIX_ROW_PADDING); + const int64_t J_max = ggml_cuda_mmq_get_J_max(type_x, fallback, cc, src1_ncols); + const size_t nbytes_src1_q8_1 = src1_ncols * ne10_padded * sizeof(block_q8_1_mmq) / QK8_1_MMQ + + J_max * sizeof(block_q8_1_mmq); + + ggml_cuda_pool_alloc src1_q8_1(pool_slice, nbytes_src1_q8_1); + + // zero the J tail so the kernel's speculative block reads beyond the padded rows are benign + fprintf(stderr, "PR104 mmq split: dev=%d type=%s ncols=%d ne10=%d nbytes=%zu ptr=%p J=%d fallback=%d\n", + ggml_cuda_get_device(), ggml_type_name(type_x), (int) src1_ncols, (int) src1->ne[0], nbytes_src1_q8_1, (void *) src1_q8_1.get(), (int) J_max, (int) fallback); + CUDA_CHECK(cudaMemsetAsync(src1_q8_1.get(), 0, nbytes_src1_q8_1, stream)); + + quantize_mmq_q8_1_cuda(src1_ddf_i, nullptr, src1_q8_1.get(), type_x, + src1->ne[0], src1->ne[0], 1, 1, ne10_padded, src1_ncols, 1, 1, stream); + CUDA_CHECK(cudaGetLastError()); + + // stride into this device's real dst region: main device writes global rows in place, + // other devices write a token-major partial with pitch row_diff + const int id = ggml_cuda_get_device(); + const int64_t s1 = id == ctx.device ? dst->nb[1] / ggml_type_size(dst->type) : row_diff; + + const mmq_args args = { + src0_dd_i, type_x, (const int *) src1_q8_1.get(), nullptr, nullptr, dst_dd_i, nullptr, + ne00, row_diff, src1_ncols, src0->nb[1] / ggml_type_size(src0->type), src1_ncols, s1, + 1, 1, 0, 0, 0, + 1, 1, 0, 0, 0, + 1}; + + ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); + + GGML_UNUSED_VARS(src1, dst, src1_padded_row_size); +} + bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts) { #ifdef GGML_CUDA_FORCE_CUBLAS return false; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index b4a747720f77..a2002414ec5c 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1596,3 +1596,10 @@ void ggml_cuda_mul_mat_q( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst); bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts); + +// PR104.2: per-op MUL_MAT for CUDA row-split buffers (see mmq.cu) +void ggml_cuda_op_mul_mat_q( + ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, + const char * src0_dd_i, const float * src1_ddf_i, float * dst_dd_i, + const int64_t row_low, const int64_t row_high, const int64_t src1_ncols, + const int64_t src1_padded_row_size, cudaStream_t stream, const ggml_type type_slice, ggml_cuda_pool & pool_slice); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index f65e0fbcd7b0..d1242d8ef66d 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1476,7 +1476,9 @@ void ggml_cuda_op_mul_mat_vec_q( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i, const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols, - const int64_t src1_padded_row_size, cudaStream_t stream) { + const int64_t src1_padded_row_size, cudaStream_t stream, const ggml_type type_slice, ggml_cuda_pool & pool_slice) { + + const ggml_type type_x = type_slice == GGML_TYPE_COUNT ? src0->type : type_slice; const int64_t ne00 = src0->ne[0]; const int64_t row_diff = row_high - row_low; @@ -1492,13 +1494,14 @@ void ggml_cuda_op_mul_mat_vec_q( // nrows_dst == nrows of the matrix that the kernel writes into const int64_t nrows_dst = id == ctx.device ? ne0 : row_diff; - const int stride_row_x = ne00 / ggml_blck_size(src0->type); + const int stride_row_x = ne00 / ggml_blck_size(type_x); const int stride_col_y = src1_padded_row_size / QK8_1; ggml_cuda_mm_fusion_args_device fusion_local{}; mul_mat_vec_q_switch_type( - src0_dd_i, src0->type, src1_ddq_i, nullptr, fusion_local, dst_dd_i, ne00, row_diff, src1_ncols, stride_row_x, stride_col_y, nrows_dst, + src0_dd_i, type_x, src1_ddq_i, nullptr, fusion_local, dst_dd_i, ne00, row_diff, src1_ncols, stride_row_x, stride_col_y, nrows_dst, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, stream); GGML_UNUSED_VARS(src1, dst, src1_ddf_i, src1_ncols, src1_padded_row_size); + GGML_UNUSED(pool_slice); } diff --git a/ggml/src/ggml-cuda/mmvq.cuh b/ggml/src/ggml-cuda/mmvq.cuh index 5605bf7a4e60..5cc1854cecda 100644 --- a/ggml/src/ggml-cuda/mmvq.cuh +++ b/ggml/src/ggml-cuda/mmvq.cuh @@ -15,4 +15,4 @@ void ggml_cuda_op_mul_mat_vec_q( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i, const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols, - const int64_t src1_padded_row_size, cudaStream_t stream); + const int64_t src1_padded_row_size, cudaStream_t stream, const ggml_type type_slice, ggml_cuda_pool & pool_slice); From 3915308fcb8541ef958ad76559ed52e1810a3364 Mon Sep 17 00:00:00 2001 From: Hydra Engineering Date: Fri, 11 Sep 2026 01:05:47 +0700 Subject: [PATCH 8/8] docs: PR104.2 per-op row-split MUL_MAT lands; 27B mixed-quant fits and runs correct --- docs/arms/pr104-mixed-quant-row-sharding.md | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/docs/arms/pr104-mixed-quant-row-sharding.md b/docs/arms/pr104-mixed-quant-row-sharding.md index 74e588b3eba6..efee99495d13 100644 --- a/docs/arms/pr104-mixed-quant-row-sharding.md +++ b/docs/arms/pr104-mixed-quant-row-sharding.md @@ -361,3 +361,21 @@ above), `nvidia-smi` before/after 1/15712→15847 MiB. Commit on 4. Keep: capture gate, fusion gates, gather for non-MUL_MAT split sources. Then 27B mixed-quant boot + `LLAMA_ARG_SPLIT_ROW_QUANT=q4_k,q6_k` MTP production run, `test-suite.sh `, decode t/s vs ≥45 bar, PR104.0 projection comparison, fork push. **Commit:** fix 1-3 + validation is `40ea56978 cuda: harden split-buffer compute path` on `fork/pr104-mixed-quant-arm` (PR #113). Production pod lifecycle protocol followed this session: `podman pod stop pod_llama-baseline` before test boots (VRAM drained 15847/11911 → 1/1 MiB), `podman pod start` restored — health 200 on :18081, VRAM back to 15847/11911 MiB. + +## PR104.2 Progress 3 — per-op row-split MUL_MAT lands; 27B mixed-quant fits and runs correct at 12.2 t/s (perf gap open) — 2026-09-11 (session 3) + +**What shipped (commit `a38976d98`, branch `fork/pr104-mixed-quant-arm` on ddvnguyen):** + +1. `ggml_cuda_mul_mat_split` (ggml-cuda.cu ~2262): full port of the removed per-op `ggml_cuda_op_mul_mat` — per-device weights slices, `PR104_MUL_MAT_SRC1_COL_STRIDE 128` col chunking, per-device staging of src1 f32 chunks (`cudaMemcpyPeerAsync`), per-op local quantization, event-based cross-device sync (`events[dev][is]`), token-major dst partials stitched via `ggml_cuda_Memcpy2DPeerAsync` (`cudaMemcpy3DPeerAsync`, vmm-safe). No contiguous gather or weight copies in this path. +2. Per-op `ggml_cuda_op_mul_mat_q` (mmq.cu) and slice-type support added to per-op `ggml_cuda_op_mul_mat_vec_q` (mmvq.cu): both switch on the **per-device slice type** (`extra->slice_type[id]`, mixed-quant q4_K/q6_K slices), take an explicit per-device pool (`ctx.pool(id)`) — the naive `ctx.pool()` was the main-device pool crash reported as memset "invalid argument" at mmq.cu:291. +3. `compute_forward`: MUL_MAT with split src0 now skips the per-op contiguous gather; `ggml_cuda_mul_mat` dispatches to the split path first. +4. `init_tensor`/`set_tensor`: contiguous copies gated by `GGML_SPLIT_CONTIG_MAX_BYTES` (default 16 MiB). The earlier default (256 MiB/tensor → 256MiB×61 weights ≈ +12 GB) was the first 27B OOM; with 16 MiB, only tiny f32 norm/bias tensors get gathers. + +**Empirical results:** + +- 9B mixed-quant (`q4_K,q6_K`, `-sm row -ts 38,27`): boots, correct completions, VRAM **7373/3163 MiB** (down from 8089/3109 — big tensors no longer duplicated), decode **~30 t/s** (was 41.7 with the gather path; per-op overhead costs ~12 t/s on 9B), no CUDA errors after row-length sanity (`nlices` miscount fix). The q4_K/q6_K slices are consumed by their matching per-op kernels. +- **27B mixed-quant row-split fits and runs for the first time**: VRAM **8283/9001 MiB** (matches the PR104.0 budget table), health 200, completions correct and coherent ("The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. ..."), decode **12.16 t/s** over 192-token run. + +**Open: perf.** The per-op path decodes the 27B at 12.2 t/s (82 ms/token) vs the PR104.0 projection ≥45 (41.4 raw × MTP). Both 9B (30 vs 41.7) and 27B (12.2) show the per-op path slower than expected, so something structural dominates beyond per-launch overhead. Suspects queued (next session, one ~14-min 27B boot each): (a) graphs disabled by the split gate — measure `GGML_CUDA_GRAPHS=1` impact (the split path was designed graph-free; recorded events/streams make capture impossible as coded — the real fix may be a "capture inside the split op so the whole graph can be reused via `cudaGraph` per layer" variant); (b) per-op event `cudaStreamWaitEvent` serialization on shared events indices; (c) `cudaSetDevice` thrash (every op × 2 devices); (d) the dst staging buffer being pool-reallocated per op. The 9B (─14 t/s) and 27B (─29 t/s) deltas differ: the 27B delta includes MMQ processing on prompt chunks of 128 (the 9B's slower q4_K/q6_K slices cost less there). + +**Session state:** the pod test cycle is clean (tests then production restore). `test-suite.sh`, MTP, and the ≥45 bar are still open — this session's correctness milestone is the mixed-quant row-split boot itself.