From be28cb2ad56d06f38d9a3d85e179ece54d12de79 Mon Sep 17 00:00:00 2001
From: Donato Capitella <donato.capitella@reversec.com>
Date: Wed, 12 Nov 2025 08:28:03 +0000
Subject: [PATCH] Add HIP shuffle compatibility shim for gfx1151 builds

---
 toolboxes/Dockerfile.rocm-6.4.4-rocwmma | 19 ++++++++++++++++---
 1 file changed, 16 insertions(+), 3 deletions(-)

diff --git a/toolboxes/Dockerfile.rocm-6.4.4-rocwmma b/toolboxes/Dockerfile.rocm-6.4.4-rocwmma
index a29b034..ccc07f6 100644
--- a/toolboxes/Dockerfile.rocm-6.4.4-rocwmma
+++ b/toolboxes/Dockerfile.rocm-6.4.4-rocwmma
@@ -49,8 +49,23 @@ RUN chmod +x /opt/apply-rocwmma-fix.sh && /opt/apply-rocwmma-fix.sh /opt/llama.c
 RUN git clean -xdf \
  && git pull \
  && git submodule update --recursive \
+ && cat > /opt/llama.cpp/ggml/src/ggml-cuda/hip_shfl_fix.h <<'EOF'
+#ifndef HIP_SHFL_FIX_H
+#define HIP_SHFL_FIX_H
+#include "common.cuh"
+#ifndef __shfl_sync
+#define __shfl_sync(mask,var,srcLane,width) __shfl(var,srcLane,width)
+#endif
+#ifndef __shfl_up_sync
+#define __shfl_up_sync(mask,var,delta,width) __shfl_up(var,delta,width)
+#endif
+#ifndef __shfl_xor_sync
+#define __shfl_xor_sync(mask,var,laneMask,width) __shfl_xor(var,laneMask,width)
+#endif
+#endif
+EOF
  && bash -lc 'f=/opt/llama.cpp/ggml/src/ggml-cuda/mma.cuh; \
-              grep -q "vendors/hip.h" "$f" || sed -i '\''1i #include "vendors/hip.h"'\'' "$f"' \
+              grep -q "hip_shfl_fix.h" "$f" || sed -i '\''1i #include "hip_shfl_fix.h"'\'' "$f"' \
  && cmake -S . -B build \
       -DGGML_HIP=ON \
       -DAMDGPU_TARGETS=gfx1151 \
@@ -64,8 +79,6 @@ RUN git clean -xdf \
  && cmake --build build --config Release -- -j"$(nproc)" \
  && cmake --install build --config Release
 
-
-
 # libs
 RUN find /opt/llama.cpp/build -type f -name 'lib*.so*' -exec cp {} /usr/lib64/ \; \
  && ldconfig