[gfx1151] mul_mat_q<Q8_0,128,true> Memory Fault during large-ubatch prompt processing on b1292 — F16 works

[blackdeathdrow]

Operating System

Ubuntu

OS Details

Linux 6.17.0-1025-oem

GPU Model

gfx1151 (Radeon 8060S Graphics, Strix Halo)

llamacpp-rocm Build Version

b1292

Model Used

Voodisss/Qwen3-Reranker-0.6B-GGUF-llama_cpp:Q8_0

Does this issue also occur with Vulkan on upstream llama.cpp?

Vulkan OOMs on this iGPU (512 MiB frame buffer), cannot test

Issue Description

mul_mat_q<(ggml_type)8,128,true> crashes with Memory Fault Error (hipErrorIllegalAddress) during prompt processing at ub=16384 on gfx1151. Same bug class as upstream #24399 (CUDA/Blackwell, mul_mat_q<Q8_0,128> MMA write-back).

Command:

llama-server -m Qwen3-Reranker-0.6B.Q8_0.gguf --reranking -b 16384 -ub 16384 -c 40934 --no-mmap

Crash at n_tokens=8192, progress=0.21:

I slot print_timing: id  3 | task 1742 | prompt processing, n_tokens =   4096, progress = 0.10, t =   6.04 s / 677.83 t/s
I slot print_timing: id  3 | task 1742 | prompt processing, n_tokens =   6144, progress = 0.16, t =   9.50 s / 646.62 t/s
I slot print_timing: id  3 | task 1742 | prompt processing, n_tokens =   8192, progress = 0.21, t =  12.71 s / 644.60 t/s
I VGPU=0x5d7be22a8a40 SWq=0x7276227e6000, HWq=0x7274e0200000, id=4
I 	Dispatch Header =0xb02 (type=2, barrier=1, acquire=1, release=1), setup=3
I 	grid=[37920, 968, 1], workgroup=[32, 8, 1]
I 	private_seg_size=0, group_seg_size=57856
E :0:rocdevice.cpp :3608: Memory Fault Error [host: gtr9-pro, GPU index: 0, faulting addr: 0x726f543a1000, kernel: void mul_mat_q<(ggml_type)8, 128, true>(...)]
E ROCm error: an illegal memory access was encountered
E ggml-cuda.cu:103: ROCm error
I   current device: -1, in function ggml_backend_cuda_get_tensor_async at ggml-cuda.cu:3207
I   hipMemcpyAsync(data, (const char *) tensor->data + offset, size, hipMemcpyDeviceToHost, cuda_ctx->stream())

Backtrace:

libggml-base.so.0(ggml_abort)
libggml-hip.so.0 → mul_mat_q<Q8_0,128,true>
libllama.so.0(llama_decode)
libllama-server-impl.so(update_slots)
libllama-server-impl.so(server_queue::start_loop)
libc.so.6(__libc_start_main)

Watchdog reload after crash:

E BackendWatchdog: llama-server backend marked unavailable: backend connection failed during request
I ProcessManager: Process terminated, waiting for GPU driver cleanup...
W Router: Reloading model after backend watchdog reset

Related: #101 (gfx1151 Memory Fault on Q4_K_M Qwen3.6, b1280).

Workaround (confirmed)

F16 weights on same b1292/gfx1151/ub=16384 processed 41605 tokens with zero device exceptions. F16 uses mul_mat_f16, not mul_mat_q.

I slot print_timing: id  3 | task 2362 | prompt eval time =    1221.22 ms /   470 tokens,   384.86 t/s
I slot      release: id  3 | task 2362 | stop processing: n_tokens = 40971, truncated = 0
I slot print_timing: id  3 | task 2757 | prompt eval time =    1584.59 ms /   585 tokens,   369.18 t/s
I slot      release: id  3 | task 2757 | stop processing: n_tokens = 41605, truncated = 0

Multiple subsequent /v1/reranking requests returned 200 with no errors.

Additional Information

device_info: ROCm0 : Radeon 8060S Graphics (126976 MiB, 87771 MiB free)
system_info: n_threads = 16 / 32 | ROCm : NO_VMM = 1 | PEER_MAX_BATCH_SIZE = 128
grid=[37920, 968, 1], workgroup=[32, 8, 1]

[chejh-amd]

Hi @blackdeathdrow , Might be worth double checking whether the crash really scales with ub size, for example if -ub 4096/8192 also crash and only smaller values work, that pretty much confirms it’s coming from the MMA tile write back path.