From: uvos <redacted>
Date: Tue, 12 Aug 2025 20:15:12 +0000 (+0200)
Subject: HIP: disable sync warp shuffel operators from clr amd_warp_sync_functions.h (#15273)
X-Git-Tag: upstream/0.0.6164~24
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=b0493156fa8622694f21f42460a84da3eded0bc0;p=pkg%2Fggml%2Fsources%2Fllama.cpp

HIP: disable sync warp shuffel operators from clr amd_warp_sync_functions.h (#15273)
---

diff --git a/ggml/src/ggml-cuda/fattn-wmma-f16.cu b/ggml/src/ggml-cuda/fattn-wmma-f16.cu
index fdc4d17d..6bc7943c 100644
--- a/ggml/src/ggml-cuda/fattn-wmma-f16.cu
+++ b/ggml/src/ggml-cuda/fattn-wmma-f16.cu
@@ -15,7 +15,6 @@ namespace wmma = mtmusa::wmma;
 namespace wmma = nvcuda::wmma;
 #endif // GGML_USE_MUSA
 #elif defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)
-#undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers
 #include <rocwmma/rocwmma.hpp>
 namespace wmma = rocwmma;
 #endif // !defined(GGML_USE_HIP)
diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h
index c31f3192..96f8bc75 100644
--- a/ggml/src/ggml-cuda/vendors/hip.h
+++ b/ggml/src/ggml-cuda/vendors/hip.h
@@ -1,6 +1,6 @@
 #pragma once
 
-#define HIP_ENABLE_WARP_SYNC_BUILTINS 1
+#define HIP_DISABLE_WARP_SYNC_BUILTINS 1
 #include <hip/hip_runtime.h>
 #include <hipblas/hipblas.h>
 #include <hip/hip_fp16.h>