Oh thank you for the map Mario, just updated the bug ** Summary changed:
- libggml0-backend-hip: NaN / garbage tokens on gfx1152 (Strix Halo) due to mmvq nwarps host/device mismatch under gfx11-generic + libggml0-backend-hip: NaN / garbage tokens on gfx1152 (Kracken) due to mmvq nwarps host/device mismatch under gfx11-generic ** Description changed: [Impact] - On AMD Strix Halo APUs (Radeon 860M, gfx1152), any llama.cpp / ggml + On AMD Kracken APUs (Radeon 860M, gfx1152), any llama.cpp / ggml workload that uses the HIP backend and hits the mul_mat_vec_q (mmvq) kernel with quantized weights produces NaN outputs. In practice this means llama-cli emits an endless stream of "?????" / garbage tokens instead of coherent text. - The issue is not specific to Strix Halo hardware — it is triggered + The issue is not specific to Kracken hardware — it is triggered whenever the HIP runtime loads the gfx11-generic code object for an RDNA3.5 device. libggml-hip.so in Ubuntu is built with --offload-arch=gfx11-generic (no native gfx1150/51/52/53 code objects), so every RDNA3.5 part on Ubuntu is affected. [Test Case] On a gfx1152 machine (e.g. Ryzen AI Max+ 395 / Radeon 860M): $ llama-cli -m <any Q4_0 gguf> -p "hello" Expected: coherent tokens. Observed: NaN logits, output is "?????..." forever. [Root Cause] src/ggml-cuda/mmvq.cu has two get_device_table_id(): * __device__ variant, selected at compile time from RDNA3_0/RDNA3_5 macros defined in vendors/hip.h. With --offload-arch=gfx11-generic only __GFX11__ is defined; __gfx1150/51/52/53__ are not. So RDNA3_5 is NOT defined, RDNA3_0 IS defined, and the device picks MMVQ_PARAMETERS_RDNA3_0 → calc_nwarps() = 8 for common quant types at ncols_dst=1. The kernel is compiled with __launch_bounds__(8*32, 1) and sizes __shared__ tmp_shared for an 8-way inter-warp reduction. * __host__ variant, called at launch time with cc = 0x1152. GGML_CUDA_CC_IS_RDNA3_5(cc) is true → MMVQ_PARAMETERS_RDNA2 → calc_nwarps() = 1 → block_dims.y = 1. The launch only populates warp 0 of tmp_shared, but the kernel's reduction loop reads nwarps-1 = 7 warps of uninitialized LDS. That produces NaN, which propagates through the model. The mismatch also exists in the reverse direction: a build that ships only gfx11-generic and runs on true RDNA3.0 hardware would hit the same class of bug from the opposite side. [Fix] Instead of recomputing nwarps on the host and hoping it agrees with the device's compile-time choice, query the compiled kernel's actual __launch_bounds__ via cudaFuncGetAttributes / hipFuncGetAttributes and derive nwarps from attr.maxThreadsPerBlock. Host then always matches whatever code object HIP loaded, regardless of native vs. generic fatbin selection. Applied in Ubuntu ggml 0.22.0-1ubuntu1 as: debian/patches/generics-fix-mmvq-nwarps-host-device-mismatch.patch [Upstream] Forwarded to upstream as a bug report (no PR — patch is documented inline in the issue thread): https://github.com/ggml-org/llama.cpp/issues/25620 -- You received this bug notification because you are a member of Ubuntu Bugs, which is subscribed to Ubuntu. https://bugs.launchpad.net/bugs/2166361 Title: libggml0-backend-hip: NaN / garbage tokens on gfx1152 (Kracken) due to mmvq nwarps host/device mismatch under gfx11-generic To manage notifications about this bug go to: https://bugs.launchpad.net/ubuntu/+source/ggml/+bug/2166361/+subscriptions -- ubuntu-bugs mailing list [email protected] https://lists.ubuntu.com/mailman/listinfo/ubuntu-bugs
