mirror/LocalAI
1
0
Fork 0
mirror of https://github.com/mudler/LocalAI.git synced 2026-07-30 09:57:57 -04:00
Code Issues Packages Projects Releases Wiki Activity

fix(buun-llama-cpp): pass WARP_SIZE to argmax __shfl_xor_sync calls

Browse Source 
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 <mudler@localai.io>
This commit is contained in:
Ettore Di Giacinto
2026-04-24 16:29:29 +00:00
committed by localai-org-maint-bot
parent a1d09a3919
commit 5b126f49e1
1 changed files with 32 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

32
backend/cpp/buun-llama-cpp/patches/0002-argmax-shfl-xor-sync-add-width.patch Normal file
View File

@@ -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;