Skip to content

Stage HIP device weight uploads through a bounded host buffer - #1070

Merged
0xShug0 merged 1 commit into
0xShug0:mainfrom
Ramon-Balaguer:fix/hip-staged-weight-upload
Oct 25, 2026
Merged

0xShug0 merged 1 commit into
0xShug0:mainfrom
Ramon-Balaguer:fix/hip-staged-weight-upload

Conversation

@Ramon-Balaguer

@Ramon-Balaguer Ramon-Balaguer commented Oct 24, 2026 •

Copy link
Copy Markdown
Contributor

Summary

On a ROCm iGPU without XNACK (Strix Halo, gfx1151), Fish Audio Drama 3 never finishes loading: the backbone weight upload stops at about 30% and the server never becomes ready. This PR makes set_tensor_bytes fill HIP device buffers through a bounded 64 MiB host staging buffer instead of passing the memory-mapped GGUF straight to ggml_backend_tensor_set. Host buffers and the other backends (CUDA, Vulkan, Metal) keep the direct copy, so their code path does not change.

With the change, the same model loads in 15-18 s, and file-backed RSS stays below ~350 MB instead of growing to ~14 GB.

Root cause

GgufTensorSource::set_backend_tensor passes set_tensor_bytes a pointer into the MAP_PRIVATE mapping of the GGUF. Each device tensor is therefore uploaded with a single hipMemcpyAsync whose source is pageable, file-backed memory, often hundreds of MB long. On this device the HIP runtime reports HMM support: 1, XNACK: 0, Direct host access: 0. Evidence from the stalled load:

  • AMD_LOG_LEVEL=4 shows one call per tensor, sourced from the GGUF mapping: hipMemcpyAsync(dst, 0x7197ebfc7e20, 290635776, hipMemcpyHostToDevice, ...).
  • In /proc/<pid>/status, RssFile grows with the upload and is never released (up to ~13.7 GB), even though discard_range() calls madvise(MADV_DONTNEED) after each tensor. RssAnon stays below 300 MB.
  • In /proc/<pid>/task/*/wchan, the loader thread blocks in svm_range_set_attr (KFD SVM, ioctl). Once RssFile nears the container's 12 GB memory limit, disk reads drop to 0 MB/s and the thread spins in user space. upload_progress stays at 0.30-0.31 for minutes.

Staging the data through a reusable heap buffer turns this into 64 MiB copies from that buffer. RssFile stays flat and the load completes.

The stall does not come from the recent Drama 3 changes. It reproduces with a local build of c8365f39 and with 2f6c70bf, a7f18608 and dd92e32b, with GGML_HIP_HIPBLASLT both ON and OFF, and with both the @builtin and @gguf model contracts.

Change

  • core::is_hip_buffer(ggml_backend_buffer_t) (new, backend.h/backend.cpp): true when the buffer's device belongs to the HIP registry. It uses the same reg_name_matches lookup that backend_type() uses for backend handles.
  • set_tensor_bytes() (tensor_source.cpp): for a HIP device buffer (the tensor's own buffer, or its view source's), the bytes are copied in chunks of at most 64 MiB into a local staging vector, and each chunk is uploaded with ggml_backend_tensor_set(tensor, staging, offset, count). Any other buffer uses the previous single ggml_backend_tensor_set call. The staging vector is freed when the call returns.

llama.cpp's loader follows the same pattern: it fills device buffers from staging buffers rather than straight from the mapped file.

Validation

Hardware and software: AMD Radeon 8060S (Strix Halo, gfx1151) with a 96 GiB VRAM carve-out (UMA), ROCm 10.0.0 (HIP 7.15.26333, AMD clang 23.0.0git), running in a Debian 13 LXC with a 12 GB memory limit.

Builds:

# default HIP settings (hipBLASLt ON)
scripts/build_linux.sh --backend hip --gpu-targets gfx1151 --model-set custom --models fish_drama3 \
  --target audiocpp_server --jobs 4 --build-type Release --build-dir build/hip-pr-default
# rocBLAS variant: same, then
cmake -S . -B build/hip-rocblas -DGGML_HIP_HIPBLASLT=OFF && cmake --build build/hip-rocblas --target audiocpp_server

Model: drama-3-backbone-4bit-experts-Q4_K_S.gguf (Jmica/Drama-3-GGUF). Server config: backend: hip, cpu_expert_layers: 0, context_length: 4096, in offline and streaming modes. Load check:

audiocpp_server --config drama3-server.json --host 127.0.0.1 --port 10099 --log --log-file load.log --load-progress 1
Build Before After (this PR on dd92e32b)
hipBLASLt ON (default) a7f18608: not ready after 5 min; 2f6c70bf: not ready after 6 min 40 s (upload_progress stuck at ~0.31) ready in 18 s, RssFile ≤ 342 MB
hipBLASLt OFF dd92e32b: stuck at 0.31 after 90 s ready in 15-18 s, RssFile ≤ 310 MB

Requests after the fix used Spanish sentences with a cloned voice, sent sequentially to one session, with fish_drama3.device_sampling=false (see below):

Mode First audio RTF
offline (response_format: wav) the whole clip 0.88-0.95
streaming (stream_format: sse, response_format: pcm) 0.6-1.1 s 0.83-0.96

The results are the same with both builds. Whisper transcripts of the generated audio match the input text. Peak VRAM with the model loaded is 21.8 GiB, the same as an earlier working build.

Regression check on the staged path: an offline request with seed: 42, run twice with 64 MiB staging chunks and twice with a 1 MiB build of the same code, gave byte-identical WAVs in all four runs (MD5 188e6fccf8881a6127ff7e92f6098906). The uploaded weights therefore do not depend on how the copy is split.

Not tested: CUDA, Vulkan and Metal. Their uploads take the same code path as before.

Related, not addressed here

On the same HIP setup, fish_drama3.device_sampling=true (the default since #1063) aborts on the first request:

ggml_cuda_compute_forward: ARGSORT failed
ROCm error: invalid configuration argument
  current device: 0, in function ggml_cuda_compute_forward at external/ggml/src/ggml-cuda/ggml-cuda.cu:3372

fish_drama3.device_sampling=false works, and all the numbers above use it. I can open a separate issue with logs.

GGUF tensor uploads into HIP device buffers now copy through a 64 MiB
host staging buffer instead of handing the memory-mapped model file
straight to ggml_backend_tensor_set. On ROCm iGPUs without XNACK
(Strix Halo, gfx1151) the direct copies left the file pages registered
with the GPU: they were never released and the load stalled at about
30% of the weights once they approached the container memory limit.

Host buffers and the other backends keep the direct copy. The new
core::is_hip_buffer() identifies the owning registry the same way
backend_type() does for backend handles.
@0xShug0
0xShug0 merged commit 2f6b42d into 0xShug0:main Oct 25, 2026
12 checks passed
@0xShug0

0xShug0 commented Oct 25, 2026

Copy link
Copy Markdown
Owner

@Ramon-Balaguer Thanks! PR merged.

cunba-ai pushed a commit to cunba-ai/audio.cpp that referenced this pull request Oct 28, 2026
…d staging) + Earshot VAD + BreezeTTS CUDA opts + Echo-TTS CFG + /health version

Upstream commits 899d84a..71f42b8 (22):
- fix HIP VMM access updates for growing pools (0xShug0#1059)
- gguf: numeric array metadata in one bulk fread (0xShug0#1072)
- Echo-TTS optional Joint-CFG (0xShug0#773); Sway time scheduling later
- validated HIP VMM defaults in Arch AMD installer (0xShug0#1080)
- HIP argsort extended to large rows (0xShug0#1076)
- Irodori-TTS MeanFlow support (0xShug0#1083)
- assets: direct uploads preserved in weight-norm tensor views (0xShug0#1081)
- HIP device weight uploads staged through bounded host buffer (0xShug0#1070)
- empty/null speech voice values treated as omitted (0xShug0#1086)
- redundant VoxCPM2 KV cache staging/reset avoided (0xShug0#1085)
- Drama3 device sampling default off (0xShug0#1087)
- Studio LoRA controls + shared adapter groups (0xShug0#1088)
- Arena styling alignment (0xShug0#1089)
- Earshot VAD native C++ port (0xShug0#1095)
- ASR VAD chunking loads any built-in VAD family (0xShug0#1096)
- /health reports build version + git commit (0xShug0#1097)
- BreezeTTS opt-in fused depth graph CUDA (0xShug0#1090); bf16 rounding fusion in rms_norm/add/SwiGLU (0xShug0#1091); batch-2 CFG decode (0xShug0#1092); option names normalized
- ci: best-effort HIP test release binaries (0xShug0#1100)
- Echo-TTS Sway time scheduling (0xShug0#1105)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants