From 5b126f49e17b02c88d1c52bfe91d92cdd9a83009 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 24 Apr 2026 16:29:29 +0000 Subject: [PATCH] fix(buun-llama-cpp): pass WARP_SIZE to argmax __shfl_xor_sync calls Two call sites in ggml/src/ggml-cuda/argmax.cu (the top-K intra-warp merge added by buun) use the 3-arg CUDA form __shfl_xor_sync(mask, var, laneMask), omitting the optional width parameter. The hipification shim at ggml/src/ggml-cuda/vendors/hip.h:33 is a function-like macro that requires all four arguments, so hipcc fails with: argmax.cu:265: too few arguments provided to function-like macro invocation note: macro '__shfl_xor_sync' defined here: #define __shfl_xor_sync(mask, var, laneMask, width) \ __shfl_xor(var, laneMask, width) Every other call in the same file already passes WARP_SIZE explicitly; aligning these two with that convention fixes the hipblas build without changing CUDA codegen (warpSize is the CUDA default). Assisted-by: Claude:claude-opus-4-7 Signed-off-by: Ettore Di Giacinto --- .../0002-argmax-shfl-xor-sync-add-width.patch | 32 +++++++++++++++++++ 1 file changed, 32 insertions(+) create mode 100644 backend/cpp/buun-llama-cpp/patches/0002-argmax-shfl-xor-sync-add-width.patch diff --git a/backend/cpp/buun-llama-cpp/patches/0002-argmax-shfl-xor-sync-add-width.patch b/backend/cpp/buun-llama-cpp/patches/0002-argmax-shfl-xor-sync-add-width.patch new file mode 100644 index 000000000..f1df67257 --- /dev/null +++ b/backend/cpp/buun-llama-cpp/patches/0002-argmax-shfl-xor-sync-add-width.patch @@ -0,0 +1,32 @@ +Subject: [PATCH] ggml-cuda/argmax: pass WARP_SIZE to the top-K __shfl_xor_sync calls + +Two __shfl_xor_sync calls in the top-K intra-warp merge drop the `width` +argument and rely on the CUDA default (warpSize). Every other call in +the same file already passes WARP_SIZE explicitly, and the HIP/ROCm +compatibility shim at ggml/src/ggml-cuda/vendors/hip.h:33 is a 4-arg +function-like macro — so the 3-arg form fails to preprocess when +building with hipcc against ROCm: + + argmax.cu:265: error: too few arguments provided to function-like + macro invocation + note: macro '__shfl_xor_sync' defined here: + #define __shfl_xor_sync(mask, var, laneMask, width) \ + __shfl_xor(var, laneMask, width) + +Align the two call sites with the rest of the file by passing WARP_SIZE +explicitly. On CUDA the generated code is unchanged (warpSize is the +default); on HIP it now matches the macro's arity. + +--- a/ggml/src/ggml-cuda/argmax.cu ++++ b/ggml/src/ggml-cuda/argmax.cu +@@ -262,8 +262,8 @@ + // Each step: lane gets partner's min element, if it beats our min, replace and re-heapify + for (int offset = WARP_SIZE / 2; offset > 0; offset >>= 1) { + for (int i = 0; i < K; i++) { +- float partner_val = __shfl_xor_sync(0xFFFFFFFF, heap_val[i], offset); +- int partner_idx = __shfl_xor_sync(0xFFFFFFFF, heap_idx[i], offset); ++ float partner_val = __shfl_xor_sync(0xFFFFFFFF, heap_val[i], offset, WARP_SIZE); ++ int partner_idx = __shfl_xor_sync(0xFFFFFFFF, heap_idx[i], offset, WARP_SIZE); + if (partner_val > heap_val[0]) { + heap_val[0] = partner_val; + heap_idx[0] = partner_idx;