Add HIP shuffle compatibility shim for gfx1151 builds
This commit is contained in:
@@ -49,8 +49,23 @@ RUN chmod +x /opt/apply-rocwmma-fix.sh && /opt/apply-rocwmma-fix.sh /opt/llama.c
|
|||||||
RUN git clean -xdf \
|
RUN git clean -xdf \
|
||||||
&& git pull \
|
&& git pull \
|
||||||
&& git submodule update --recursive \
|
&& 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; \
|
&& 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 \
|
&& cmake -S . -B build \
|
||||||
-DGGML_HIP=ON \
|
-DGGML_HIP=ON \
|
||||||
-DAMDGPU_TARGETS=gfx1151 \
|
-DAMDGPU_TARGETS=gfx1151 \
|
||||||
@@ -64,8 +79,6 @@ RUN git clean -xdf \
|
|||||||
&& cmake --build build --config Release -- -j"$(nproc)" \
|
&& cmake --build build --config Release -- -j"$(nproc)" \
|
||||||
&& cmake --install build --config Release
|
&& cmake --install build --config Release
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
# libs
|
# libs
|
||||||
RUN find /opt/llama.cpp/build -type f -name 'lib*.so*' -exec cp {} /usr/lib64/ \; \
|
RUN find /opt/llama.cpp/build -type f -name 'lib*.so*' -exec cp {} /usr/lib64/ \; \
|
||||||
&& ldconfig
|
&& ldconfig
|
||||||
|
|||||||
Reference in New Issue
Block a user