Force-include HIP shuffle shim to fix missing __shfl_sync on gfx1151 builds
This commit is contained in:
@@ -46,7 +46,8 @@ COPY ./apply-rocwmma-fix.sh /opt/apply-rocwmma-fix.sh
|
|||||||
RUN chmod +x /opt/apply-rocwmma-fix.sh && /opt/apply-rocwmma-fix.sh /opt/llama.cpp
|
RUN chmod +x /opt/apply-rocwmma-fix.sh && /opt/apply-rocwmma-fix.sh /opt/llama.cpp
|
||||||
|
|
||||||
# build
|
# build
|
||||||
RUN git clean -xdf \
|
RUN set -euo pipefail \
|
||||||
|
&& 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'
|
&& cat > /opt/llama.cpp/ggml/src/ggml-cuda/hip_shfl_fix.h <<'EOF'
|
||||||
@@ -65,8 +66,8 @@ RUN git clean -xdf \
|
|||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
EOF
|
EOF
|
||||||
&& f=/opt/llama.cpp/ggml/src/ggml-cuda/mma.cuh \
|
# remove any old inline hack you had in mma.cuh (safe if absent)
|
||||||
&& grep -q 'hip_shfl_fix.h' "$f" || sed -i '1i #include "hip_shfl_fix.h"' "$f" \
|
&& sed -i '/HIP_HAS_SHFL_SYNC_FUNCS/,+20d' /opt/llama.cpp/ggml/src/ggml-cuda/mma.cuh || true \
|
||||||
&& cmake -S . -B build \
|
&& cmake -S . -B build \
|
||||||
-DGGML_HIP=ON \
|
-DGGML_HIP=ON \
|
||||||
-DAMDGPU_TARGETS=gfx1151 \
|
-DAMDGPU_TARGETS=gfx1151 \
|
||||||
@@ -76,7 +77,7 @@ EOF
|
|||||||
-DROCM_PATH=/opt/rocm \
|
-DROCM_PATH=/opt/rocm \
|
||||||
-DHIP_PATH=/opt/rocm \
|
-DHIP_PATH=/opt/rocm \
|
||||||
-DHIP_PLATFORM=amd \
|
-DHIP_PLATFORM=amd \
|
||||||
-DCMAKE_HIP_FLAGS="--rocm-path=/opt/rocm" \
|
-DCMAKE_HIP_FLAGS="--rocm-path=/opt/rocm -include /opt/llama.cpp/ggml/src/ggml-cuda/hip_shfl_fix.h -Wno-macro-redefined" \
|
||||||
&& cmake --build build --config Release -- -j"$(nproc)" \
|
&& cmake --build build --config Release -- -j"$(nproc)" \
|
||||||
&& cmake --install build --config Release
|
&& cmake --install build --config Release
|
||||||
|
|
||||||
|
|||||||
Reference in New Issue
Block a user