try AMD fix

Author: JohannesGaessler

ggml/src/ggml-cuda/common.cuh

@@ -420,12 +420,9 @@ static __device__ __forceinline__ half2 warp_reduce_sum(half2 a) {
 
 template<int width = WARP_SIZE>
 static __device__ __forceinline__ int warp_reduce_all(int x) {
-#ifndef GGML_USE_HIP
-    if (width == WARP_SIZE) {
+    if (width == ggml_cuda_get_physical_warp_size()) {
         return __all_sync(0xffffffff, x);
-    } else
-#endif // GGML_USE_HIP
-    {
+    } else {
 #pragma unroll
         for (int offset = width/2; offset > 0; offset >>= 1) {
             x = __shfl_xor_sync(0xffffffff, x, offset, width) && x;
@@ -436,12 +433,9 @@ static __device__ __forceinline__ int warp_reduce_all(int x) {
 
 template<int width = WARP_SIZE>
 static __device__ __forceinline__ int warp_reduce_any(int x) {
-#ifndef GGML_USE_HIP
-    if (width == WARP_SIZE) {
+    if (width == ggml_cuda_get_physical_warp_size()) {
         return __any_sync(0xffffffff, x);
-    } else
-#endif // GGML_USE_HIP
-    {
+    } else {
 #pragma unroll
         for (int offset = width/2; offset > 0; offset >>= 1) {
             x = __shfl_xor_sync(0xffffffff, x, offset, width) || x;

ggml/src/ggml-cuda/mmq.cu

@@ -41,7 +41,7 @@ static __global__ void mmq_ids_helper(
                 ids_dst_shared[it_compact]  = it*n_expert_used + iex_used;
             }
 
-            if (warp_reduce_any(iex_used != -1)) {
+            if (warp_reduce_any<warp_size>(iex_used != -1)) {
                 it_compact++;
             }
         }
@@ -80,7 +80,7 @@ static __global__ void mmq_ids_helper(
             it_compact += __shfl_sync(0xFFFFFFFF, it_compact_add_lower + it_compact_add_self, warp_size - 1, warp_size);
         }
     }
-    nex_prev = warp_reduce_sum(nex_prev);
+    nex_prev = warp_reduce_sum<warp_size>(nex_prev);
 
     for (int it = threadIdx.x; it < it_compact; it += warp_size) {
         ids_src1[nex_prev + it] = ids_src1_shared[it];

ggml/src/ggml-cuda/vendors/hip.h

@@ -24,6 +24,8 @@
 #define __shfl_sync(mask, var, laneMask, width) __shfl(var, laneMask, width)
 #define __shfl_up_sync(mask, var, laneMask, width) __shfl_up(var, laneMask, width)
 #define __shfl_xor_sync(mask, var, laneMask, width) __shfl_xor(var, laneMask, width)
+#define __all_sync(mask, var) __all(var)
+#define __any_sync(mask, var) __any(var)
 #define cublasCreate hipblasCreate
 #define cublasDestroy hipblasDestroy
 #define cublasGemmEx hipblasGemmEx