[Bug] ROCm dual-device: streaming diffusion buffer init fails with "unspecified launch failure / current device: -1" when a gfx1150 iGPU is visible alongside gfx1100 dGPU (MiniMax-H3 ref2va)

Git commit

95b3090 (master-975 prebuilt, official Windows ROCm bundle; prebuilt carries no source, so exact ggml submodule state unknown — crash path: ggml/src/ggml-cuda/ggml-cuda.cu:817)

Operating System & Version

Windows 11 (AMD platform)

GGML backends

HIP

Command-line arguments used

Failing case (all three of these variants fail; shown: variant C):

sd-cli -M vid_gen \
  --diffusion-model minimax_h3_ref2va_pruned-Q4_K_M.gguf \
  --vae minimax_h3_video_vae_fp16.safetensors \
  --llm qwen3vl_32b_minimax_h3-Q4_K_M.gguf \
  --cfg-scale 1 -W 480 -H 864 --fps 24 --video-frames 13 \
  --flash-attn off,diffusion=on \
  --conv-direct auto,diffusion=on,controlnet=on,vae=on \
  --rng cpu --flow-shift 12 --steps 4 --sampling-method euler --scheduler discrete \
  --log-level debug -t 16 \
  --backend "llm=ROCm0,vae=ROCm1,diffusion=ROCm1" \
  --params-backend "diffusion=cpu,te=cpu" \
  --max-vram "ROCm0=12" \
  --ref-image <one image> -p "a cat sitting on a table" -o out.mp4

Minimal repro variant (recommended to reproduce — variant B, the iGPU is NEVER used for compute):

--backend "llm=ROCm1,vae=ROCm1,diffusion=ROCm1"
--params-backend "diffusion=cpu,te=cpu"
--max-vram "ROCm0=12"
(HIP_VISIBLE_DEVICES unset, so both devices are visible; iGPU is never initialized/used)

Steps to reproduce

Machine: 64 GB system RAM, AMD iGPU Radeon 890M (gfx1150, BIOS dedicated allocation = minimum 0.5 GB, shared pool enabled) + discrete RX 7900 XTX (gfx1100).
Ensure both devices are visible to HIP (no HIP_VISIBLE_DEVICES filter). The iGPU is enumerated as ROCm0 with ~43.7 GB of shared VRAM.
Run any MiniMax-H3 ref2va/video command with streaming params backend (--params-backend diffusion=cpu) and diffusion placed on the dGPU.
Crash after Qwen3VL conditioning and diffusion weight load, at the first diffusion compute-buffer init on the dGPU.

What you expected to happen

Clean generation (single-device run does: same command with only the dGPU visible via HIP_VISIBLE_DEVICES=1 completes in ~62 s and writes the video), or at minimum a clear error.

What actually happened

Deterministic crash, reproduced 3/3 back-to-back with no reboot in between (LLM stage 24.95 / 24.95 / 25.55 s), identical line every time:

[V] model manager prepared params backend buffers (10975.93 MB, 532 tensors, 1 blocks, RAM) on ROCm_Host
[E] ggml: ROCm error: unspecified launch failure
[E] ggml:   current device: -1, in function ggml_backend_cuda_buffer_init_tensor at .../ggml/src/ggml-cuda/ggml-cuda.cu:817
[E] ggml:   hipStreamSynchronize(((hipStream_t)2))

Process exits with 0xC0000409 (STATUS_STACK_BUFFER_OVERRUN, GGML abort path). No output file.

Isolation matrix (same binary, same models, same flags; only device visibility / placement / max-vram vary)

variant	devices visible	placement	--max-vram	result
C	2 (iGPU+dGPU)	llm→iGPU, vae/diff→dGPU	ROCm0=12	FAIL 3/3 (crash above)
A	1 (dGPU only, HIP=1)	all dGPU	(removed)	PASS, 62.0 s
A2	1 (dGPU only, HIP=1)	all dGPU	ROCm0=12	PASS, 71.9 s
B	2 (iGPU+dGPU)	all dGPU, iGPU backend never initialized	ROCm0=12	FAIL (same crash)

Conclusions:

--max-vram is not the variable (A2 vs B).
Using the iGPU for compute is NOT required (B): the iGPU backend is never created in B (log shows only Initializing backend: ROCm1 + CPU), yet the dGPU buffer init still fails.
The only difference between PASS and FAIL rows is whether the gfx1150 iGPU is enumerated/initialized in the process.
VAE buffer init on the dGPU succeeds in all failing runs; only the streaming diffusion compute-buffer init (ggml_backend_cuda_buffer_init_tensor with hipStreamSynchronize on a per-thread stream) fails, with current device: -1 reported by the error handler.

Not the same root cause as #2008

#2008 ("unspecified launch failure" at first compute, gfx1100, single GPU) was narrowed to pinned host memory and is avoided with GGML_CUDA_NO_PINNED=1. This report:

requires TWO HIP devices to be visible (a gfx1150 iGPU in the mix);
is independent of --max-vram value and of whether the iGPU is used;
fails at buffer init (hipStreamSynchronize), not at kernel launch after tensor load.

Possible mechanism (related upstream work)

ggml-org/llama.cpp#21140 documents the same current device: -1 family: on multi-GPU ROCm, hipGetDevice() can return an unexpected value on threads that never explicitly initialized a device context, so the early-return in ggml_cuda_set_device() skips hipSetDevice() and the next stream op fails. PR #21170 (remove that early return) is open/unmerged. If the vendored ggml in this prebuilt still has the early-return, this may be the same bug class surfacing at buffer_init_tensor when a second (gfx1150) agent is present.

Environment details

Driver: Adrenalin 32.0.31041 (ROCm 10.1 surface)
ggml init: ggml_cuda_init: found 2 ROCm devices (Total VRAM: 69256 MiB) — iGPU gfx1150 44696 MiB / dGPU gfx1100 24560 MiB
iGPU BIOS dedicated VRAM = 0.5 GB (minimum); shared memory pool enabled (HIP reports ~43.7 GB for the iGPU). Dedicated-allocation test pending (planning to raise it to 4 GB and re-run variant B; will report back).
Single-device (HIP_VISIBLE_DEVICES=1) runs work fine, including the same streaming config.

Logs

Full captured logs available (4 files: C/B failing, A/A2 passing; each includes device enumeration, the full plan fingerprint, and the 3-line error). Happy to attach or zip them here.

sdcpp-issue-roc-dual-device-buffer-init_logs.zip

[追加] Qwen Image 2.1 8B line, dual-GPU te→iGPU, new signature — same machine/95b3090/ROCm 10.1

Same prebuilt (master-975, 95b3090) + ROCm 10.1.0, machine with 890M iGPU (gfx1150) + RX 7900 XTX (gfx1100). This time the Qwen Image 2.1 image line (8B text encoder, not H3 32B), image-edit mode, 2 reference images, placement te→ROCm0 (890M), diffusion/vae→ROCm1 (7900XTX):

[D] qwen3vl executing segment 1/1: graph
[E] ggml: ROCm error: invalid configuration argument
[E] ggml: current device: 0, in function ggml_cuda_kernel_launch at common.cuh:1742
[E] ggml: hipGetLastError()

current device: 0 = the 890M is actually being used for the text encoder here (contrast: #2169's iGPU was merely visible, unused). Error captured at hipGetLastError() right after the launch — reads like a silently-failed HSA agent/context init on gfx1150, not a per-kernel launch-config violation.

Control (known-good): build 959 (d637cfb) on this same box + ROCm 10.1 ran this exact dual-GPU qwen21 use cleanly (rc=0, ~82s single image). So the regression sits in 959..975; the window's main event is the ggml submodule sync (#2147, 277 llama.cpp commits). gfx1150-relevant suspects: fa692ff (fused RMS_NORM+SCALE), 757416a (norm/rope grid-dim clamp), 8d2974e (RDNA3.5 MMQ 64-row tiles), e05d18a (new precomputed RoPE op).

Auto-fallback to single-GPU (HIP=1, te→CPU) works but CPU 8B+vision is impractically slow for this line.

Happy to run any A/B you need: 959 vs 975 side-by-side, HIP_VISIBLE_DEVICES combos, or per-flag bisect of the ggml suspects.

[10-21 update — same machine: Ryzen AI 9 HX 470 with 890M iGPU (gfx1150) + RX 7900 XTX (gfx1100), Windows 11, new AMD Adrenalin driver + reboot]

Re-tested on both prebuilts (0518b34 and 95b3090). Two findings:

1. The 8B te→iGPU coop placement now works (comment #1 signature no longer reproduces)

Qwen Image 2.1 8B with te compute on the iGPU, diffusion+VAE on the dGPU:

sd-cli -M img_gen \
  --diffusion-model qwen_image_2.1-Q8_0.gguf \
  --llm Qwen3-VL-8B-Instruct-Q8_0.gguf \
  --vae qwen_image_2.1_vae_bf16.safetensors \
  --conv-direct auto,diffusion=on,controlnet=on,vae=on --rng cpu \
  --steps 4 --cfg-scale 6.0 --sampling-method euler -H 1024 -W 1024 \
  --backend "llm=ROCm0,vae=ROCm1,diffusion=ROCm1" \
  --params-backend "diffusion=cpu,te=cpu" \
  -p "a red teacup on a wooden table" -o out.png

rc=0 on BOTH 0518b34 and 95b3090 (~32 s, 1024x1024 image fine, te 7.7GB RAM-staged with compute on ROCm0/iGPU confirmed in the log). Since the 95b3090↔0518b34 crash path has zero diff, the change is on the driver side (new Adrenalin), not sdcpp — the "invalid configuration argument / current device: 0" from the 10-19 addendum does not reproduce anymore.

2. H3 32B streaming buffer init still crashes when the iGPU is merely visible (variant B)

Minimal repro from the issue (iGPU enumerated but never used for compute, everything on ROCm1):

sd-cli -M vid_gen \
  --diffusion-model minimax_h3_ref2va_pruned-Q4_K_M.gguf \
  --vae minimax_h3_video_vae_fp16.safetensors \
  --llm qwen3vl_32b_minimax_h3-Q4_K_M.gguf \
  --cfg-scale 1 -W 480 -H 864 --fps 24 --video-frames 13 \
  --flash-attn off,diffusion=on --conv-direct auto,diffusion=on,controlnet=on,vae=on \
  --rng cpu --flow-shift 12 --steps 4 --sampling-method euler --scheduler discrete -t 16 \
  --backend "llm=ROCm1,vae=ROCm1,diffusion=ROCm1" \
  --params-backend "diffusion=cpu,te=cpu" \
  --max-vram "ROCm0=12" \
  --ref-image ref.jpg -p "a cat sitting on a table" -o out.mp4

Fails right after "loading tensors completed": [E] ggml: ROCm error: unspecified launch failure / current device: -1, in ggml_backend_cuda_buffer_init_tensor at ggml-cuda.cu:817 (hipStreamSynchronize at ggml-cuda.cu:111), process aborts with 0xC0000409. Reproduced on 0518b34, byte-identical fingerprint to the original report. So the iGPU-visible crash in the H3 32B line persists even with the new driver — only the 8B te path recovered.

Enumeration gotcha (helps repro): a stray system-level HIP_VISIBLE_DEVICES=1 env var silently filters the iGPU out of HSA enumeration (found 1 instead of found 2). Remove it in-process (Remove-Item Env:HIP_VISIBLE_DEVICES on Windows PowerShell) to get the two-device enumeration the repro needs.

Release: master-982-0518b34, asset sd-master-0518b34-bin-win-rocm-10.1.0-x64.zip (Windows ROCm 10.1.0 prebuilt)

Asset contents (verified by unzip)

amd_comgr.dll, amdhip64_7.dll, rocm_kpack.dll, sd-cli.exe, sd-server.exe, stable-diffusion.dll, start-server.bat/.sh — only the 3 runtime DLLs, no hipblas/rocblas chain, no data dirs.

Problem

stable-diffusion.dll from this build hard-imports hipblas.dll (new in the 982 build; the 975/95b3090 prebuilt did not), with the load chain:

stable-diffusion.dll → hipblas.dll → {rocblas.dll, rocsolver.dll}
                       → libhipblaslt.dll → origami.dll   (+ amdhip64_7 / 3-piece set)

(import tables verified with pefile.) So out of the box:

sd-cli.exe --version / sd-server.exe → 0xC0000135 STATUS_DLL_NOT_FOUND
Even with the DLLs resolvable, rocBLAS/hipBLASLt locate their Tensile data dirs (rocblas\library\TensileLibrary*.dat, hipblaslt\library\...) relative to the DLL's own directory (no env var exists — strings-scan of rocblas.dll / libhipblaslt.dll shows only HIPBLASLT_EXT_OP_LIBRARY_PATH, which is for extension ops, not the main libs). If the data dirs are not siblings of the DLLs, model loading silently corrupts/fails at the tail (no [E] log line).

Local install is not picked up

A locally installed ROCm 10.1 (Windows pip rocm-sdk-* wheels, DLLs in <python>\Lib\site-packages\_rocm_sdk_libraries\bin + _rocm_sdk_core) is not found by the prebuilt — Windows hard-import resolution only walks exe dir → system dirs → PATH. Two working workarounds:

Copy the hipblas chain + the two data dirs next to sd-cli.exe (full self-contained dir), or
Prepend the two site-packages bin dirs to PATH for the process (_rocm_sdk_core\bin, _rocm_sdk_libraries\bin) — the data dirs are already siblings of the DLLs there, zero copies.

Expectation

Either the win-rocm asset should bundle the hipblas chain + rocblas/library + hipblaslt/library (the 982 build needs them), or stable-diffusion.dll should not hard-import hipblas (975 behavior). As shipped, the 982 win-rocm prebuilt cannot start without the user knowing to supply ROCm library files that were not in the zip.

Repro: unzip the 982 win-rocm-10.1.0 asset on a clean Windows machine (ROCm installed via AMD's Windows pip wheels only, nothing in the exe dir besides the zip), run sd-cli.exe --version.

以上由大模型自行发现和总结的，如有错误，还请海涵。

Thanks for the detailed logs. I traced the failure to quantized weight padding initialization during diffusion weight staging, before the diffusion graph executes.

One candidate worth testing is ggml commit ee665f42 (ggml-org/llama.cpp#28782), included in #2147. It changed this exact path from cudaMemset to cudaMemsetAsync on cudaStreamPerThread, followed by synchronization. This is a candidate, not yet a confirmed cause.

Could you first rerun variant B, keeping both GPUs visible and all other settings unchanged, with:

$env:GGML_CUDA_NO_PINNED = "1"

The logs show that CPU parameters use ROCm_Host, so pinned memory has not yet been ruled out. If this still fails, the next useful comparison is a build reverting only the padding memset change.

Also, current device: -1 does not necessarily indicate an uninitialized device context: the error handler initializes the ID to -1 and ignores the return status of hipGetDevice(), which may itself fail after a GPU error.

The Qwen invalid configuration argument report may be separate; we should track it independently until we identify the failing kernel.

6. Thanks for the detailed logs. I traced the failure to quantized weight padding initialization during diffusion weight staging, before the diffusion graph executes.
One candidate worth testing is ggml commit ee665f42 (ggml-org/llama.cpp#28782), included in #2147. It changed this exact path from cudaMemset to cudaMemsetAsync on cudaStreamPerThread, followed by synchronization. This is a candidate, not yet a confirmed cause.
Could you first rerun variant B, keeping both GPUs visible and all other settings unchanged, with:
$env:GGML_CUDA_NO_PINNED = "1"
The logs show that CPU parameters use ROCm_Host, so pinned memory has not yet been ruled out. If this still fails, the next useful comparison is a build reverting only the padding memset change.
Also, current device: -1 does not necessarily indicate an uninitialized device context: the error handler initializes the ID to -1 and ignores the return status of hipGetDevice(), which may itself fail after a GPU error.
The Qwen invalid configuration argument report may be separate; we should track it independently until we identify the failing kernel.

Reran variant B as requested: 0518b34 (master-982) + ROCm 10.1 + current Adrenalin
driver, both GPUs visible (no HIP_VISIBLE_DEVICES filter), iGPU never used for
compute, all other settings unchanged; single variable = GGML_CUDA_NO_PINNED.

Leg A (baseline, NO_PINNED unset): FAIL, rc=0xC0000409 — reproduces the
original signature exactly:
[E] ggml: ROCm error: unspecified launch failure
[E] ggml: current device: -1, in function ggml_backend_cuda_buffer_init_tensor
at .../ggml-cuda.cu:817
(crash right after "staged compute params (16239.66 MB, 551 tensors, 16 blocks)
to ROCm1")
Leg B (GGML_CUDA_NO_PINNED=1): PASS, rc=0 — sampling 149.2s + VAE decode 7.2s,
video written.

So pinned host allocations (ROCm_Host) are confirmed as the trigger for the
variant-B crash while the gfx1150 iGPU is visible; the run is unblocked on the
current build with no code change. Happy to run the ee665f42 (cudaMemset ->
cudaMemsetAsync) revert-build comparison separately if you still want to pin
the exact commit. Full logs attached.

variantB_A_baseline.log
variantB_B_no_pinned_1.log
variantB_nopinned_summary.txt

Navigation Menu

Repository navigation

[Bug] ROCm dual-device: streaming diffusion buffer init fails with "unspecified launch failure / current device: -1" when a gfx1150 iGPU is visible alongside gfx1100 dGPU (MiniMax-H3 ref2va) #2169

Git commit

Operating System & Version

GGML backends

Command-line arguments used

Steps to reproduce

What you expected to happen

What actually happened

Isolation matrix (same binary, same models, same flags; only device visibility / placement / max-vram vary)

Not the same root cause as #2008

Possible mechanism (related upstream work)

Environment details

Logs

Asset contents (verified by unzip)

Problem

Local install is not picked up

Expectation

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

[Bug] ROCm dual-device: streaming diffusion buffer init fails with "unspecified launch failure / current device: -1" when a gfx1150 iGPU is visible alongside gfx1100 dGPU (MiniMax-H3 ref2va) #2169

Description

Git commit

Operating System & Version

GGML backends

Command-line arguments used

Steps to reproduce

What you expected to happen

What actually happened

Isolation matrix (same binary, same models, same flags; only device visibility / placement / max-vram vary)

Not the same root cause as #2008

Possible mechanism (related upstream work)

Environment details

Logs

Activity

farawayso commented on Oct 19, 2026

farawayso commented on Oct 21, 2026

farawayso commented on Oct 21, 2026

Asset contents (verified by unzip)

Problem

Local install is not picked up

Expectation

farawayso commented on Oct 21, 2026

leejet commented on Oct 21, 2026

farawayso commented on Oct 21, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions