Skip to content

Commit 74c7a2f

Browse files
committed
Starting from ROCm 6.5 (aka 6.4 with 7.0 preview) HIP_ENABLE_WARP_SYNC_BUILTINS has been replaced with HIP_DISABLE_WARP_SYNC_BUILTINS (https://github.com/ROCm/clr/blob/rocm-6.4.x-with-7.0-preview/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h#L30)
1 parent 4d9df1d commit 74c7a2f

File tree

1 file changed

+5
-1
lines changed

1 file changed

+5
-1
lines changed

ggml/src/ggml-cuda/fattn-wmma-f16.cu

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,11 @@ namespace wmma = mtmusa::wmma;
1515
namespace wmma = nvcuda::wmma;
1616
#endif // GGML_USE_MUSA
1717
#elif defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)
18-
#undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers
18+
#if HIP_VERSION >= 60500000
19+
#define HIP_DISABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers for ROCm 6.5+
20+
#else
21+
#undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers before ROCm 6.5
22+
#endif // HIP_VERSION >= 60500000
1923
#include <rocwmma/rocwmma.hpp>
2024
namespace wmma = rocwmma;
2125
#endif // !defined(GGML_USE_HIP)

0 commit comments

Comments
 (0)