uvos commited on
Commit
8fca6dd
·
1 Parent(s): 7b868ed

HIP: disable sync warp shuffel operators from clr amd_warp_sync_functions.h (llama/15273)

Browse files
ggml/src/ggml-cuda/fattn-wmma-f16.cu CHANGED
@@ -15,7 +15,6 @@ namespace wmma = mtmusa::wmma;
15
  namespace wmma = nvcuda::wmma;
16
  #endif // GGML_USE_MUSA
17
  #elif defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)
18
- #undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers
19
  #include <rocwmma/rocwmma.hpp>
20
  namespace wmma = rocwmma;
21
  #endif // !defined(GGML_USE_HIP)
 
15
  namespace wmma = nvcuda::wmma;
16
  #endif // GGML_USE_MUSA
17
  #elif defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)
 
18
  #include <rocwmma/rocwmma.hpp>
19
  namespace wmma = rocwmma;
20
  #endif // !defined(GGML_USE_HIP)
ggml/src/ggml-cuda/vendors/hip.h CHANGED
@@ -1,6 +1,6 @@
1
  #pragma once
2
 
3
- #define HIP_ENABLE_WARP_SYNC_BUILTINS 1
4
  #include <hip/hip_runtime.h>
5
  #include <hipblas/hipblas.h>
6
  #include <hip/hip_fp16.h>
 
1
  #pragma once
2
 
3
+ #define HIP_DISABLE_WARP_SYNC_BUILTINS 1
4
  #include <hip/hip_runtime.h>
5
  #include <hipblas/hipblas.h>
6
  #include <hip/hip_fp16.h>