Commit 278a89cb7 for llama.cpp
commit 278a89cb787c1d2b4cff9748b1df9cfd2aa69976
Author: uvos <carl@uvos.xyz>
Date: Sun Nov 15 10:58:36 2026 +0100
CUDA/HIP: Refactor the use of _sync warp primitives (#31617)
Uses the rocm-native primitives, increases the rocm version requirement to 6.3 and makes top-k selectable warp.
diff --git a/.github/workflows/build-cuda-ubuntu.yml b/.github/workflows/build-cuda-ubuntu.yml
index db97aeb15..cdbbb4767 100644
--- a/.github/workflows/build-cuda-ubuntu.yml
+++ b/.github/workflows/build-cuda-ubuntu.yml
@@ -98,7 +98,7 @@ jobs:
hip:
runs-on: ubuntu-22.04
- container: rocm/dev-ubuntu-22.04:6.1.2
+ container: rocm/dev-ubuntu-22.04:6.3.4
steps:
- name: Clone
diff --git a/ggml/src/ggml-cuda/argmax.cu b/ggml/src/ggml-cuda/argmax.cu
index 51967c667..9be901e56 100644
--- a/ggml/src/ggml-cuda/argmax.cu
+++ b/ggml/src/ggml-cuda/argmax.cu
@@ -22,8 +22,8 @@ static __global__ void argmax_f32(const float * __restrict__ x, int32_t * __rest
#pragma unroll
for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) {
- const float val = __shfl_xor_sync(0xFFFFFFFF, maxval, offset, WARP_SIZE);
- const int col = __shfl_xor_sync(0xFFFFFFFF, argmax, offset, WARP_SIZE);
+ const float val = ggml_cuda_shfl_xor(maxval, offset, WARP_SIZE);
+ const int col = ggml_cuda_shfl_xor(argmax, offset, WARP_SIZE);
if (val > maxval) {
maxval = val;
argmax = col;
@@ -51,8 +51,8 @@ static __global__ void argmax_f32(const float * __restrict__ x, int32_t * __rest
}
#pragma unroll
for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) {
- const float val = __shfl_xor_sync(0xFFFFFFFF, maxval, offset, WARP_SIZE);
- const int col = __shfl_xor_sync(0xFFFFFFFF, argmax, offset, WARP_SIZE);
+ const float val = ggml_cuda_shfl_xor(maxval, offset, WARP_SIZE);
+ const int col = ggml_cuda_shfl_xor(argmax, offset, WARP_SIZE);
if (val > maxval) {
maxval = val;
argmax = col;
diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index 414177308..bfbd6b621 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -19,7 +19,6 @@
#define GGML_COMMON_IMPL_MUSA
#endif
#endif
-#include "ggml-common.h"
#include <array>
#include <algorithm>
@@ -39,6 +38,8 @@
#include "vendors/cuda.h"
#endif // defined(GGML_USE_HIP)
+#include "ggml-common.h"
+
#define STRINGIZE_IMPL(...) #__VA_ARGS__
#define STRINGIZE(...) STRINGIZE_IMPL(__VA_ARGS__)
@@ -126,6 +127,8 @@
# define GGML_CUDA_USE_PDL
#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && (CUDART_VERSION >= 12030 || (!(defined(_MSC_VER) && !defined(__clang__)) && CUDART_VERSION >= 11080))
+constexpr ggml_cuda_thread_mask_t GGML_CUDA_FULL_THREAD_MASK = (ggml_cuda_thread_mask_t)0xffffffffffffffffULL;
+
static __device__ __forceinline__ void ggml_cuda_syncwarp() {
#ifndef GGML_USE_HIP
__syncwarp();
@@ -430,6 +433,32 @@ static constexpr __device__ int ggml_cuda_get_max_cpy_bytes() {
#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA)
}
+template <typename T>
+static __device__ __forceinline__ T ggml_cuda_shfl(T value, int src_lane, int width = ggml_cuda_get_physical_warp_size(), ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __shfl_sync(mask, value, src_lane, width);
+}
+
+template <typename T>
+static __device__ __forceinline__ T ggml_cuda_shfl_up(T value, unsigned delta, int width = ggml_cuda_get_physical_warp_size(), ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __shfl_up_sync(mask, value, delta, width);
+}
+
+template <typename T>
+static __device__ __forceinline__ T ggml_cuda_shfl_xor(T value, int lane_mask, int width = ggml_cuda_get_physical_warp_size(), ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __shfl_xor_sync(mask, value, lane_mask, width);
+}
+
+static __device__ __forceinline__ ggml_cuda_thread_mask_t ggml_cuda_ballot(int predicate, ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __ballot_sync(mask, predicate);
+}
+
+static __device__ __forceinline__ int ggml_cuda_all(int predicate, ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __all_sync(mask, predicate);
+}
+
+static __device__ __forceinline__ int ggml_cuda_any(int predicate, ggml_cuda_thread_mask_t mask = GGML_CUDA_FULL_THREAD_MASK) {
+ return __any_sync(mask, predicate);
+}
[[noreturn]]
static __device__ void no_device_code(
@@ -476,11 +505,11 @@ struct ggml_cuda_unroll<1> {
template<int width = WARP_SIZE>
static __device__ __forceinline__ int warp_reduce_sum(int x) {
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
- return __reduce_add_sync(0xffffffff, x);
+ return __reduce_add_sync(GGML_CUDA_FULL_THREAD_MASK, x);
#else
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x += __shfl_xor_sync(0xffffffff, x, offset, width);
+ x += ggml_cuda_shfl_xor(x, offset, width);
}
return x;
#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
@@ -490,7 +519,7 @@ template<int width = WARP_SIZE>
static __device__ __forceinline__ float warp_reduce_sum(float x) {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x += __shfl_xor_sync(0xffffffff, x, offset, width);
+ x += ggml_cuda_shfl_xor(x, offset, width);
}
return x;
}
@@ -499,8 +528,8 @@ template<int width = WARP_SIZE>
static __device__ __forceinline__ float2 warp_reduce_sum(float2 a) {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- a.x += __shfl_xor_sync(0xffffffff, a.x, offset, width);
- a.y += __shfl_xor_sync(0xffffffff, a.y, offset, width);
+ a.x += ggml_cuda_shfl_xor(a.x, offset, width);
+ a.y += ggml_cuda_shfl_xor(a.y, offset, width);
}
return a;
}
@@ -510,7 +539,7 @@ static __device__ __forceinline__ half2 warp_reduce_sum(half2 a) {
#ifdef FP16_AVAILABLE
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- a = __hadd2(a, __shfl_xor_sync(0xffffffff, a, offset, width));
+ a = __hadd2(a, ggml_cuda_shfl_xor(a, offset, width));
}
return a;
@@ -523,11 +552,11 @@ static __device__ __forceinline__ half2 warp_reduce_sum(half2 a) {
template<int width = WARP_SIZE>
static __device__ __forceinline__ int warp_reduce_all(int x) {
if (width == ggml_cuda_get_physical_warp_size()) {
- return __all_sync(0xffffffff, x);
+ return ggml_cuda_all(x);
} else {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x = __shfl_xor_sync(0xffffffff, x, offset, width) && x;
+ x = ggml_cuda_shfl_xor(x, offset, width) && x;
}
return x;
}
@@ -536,11 +565,11 @@ static __device__ __forceinline__ int warp_reduce_all(int x) {
template<int width = WARP_SIZE>
static __device__ __forceinline__ int warp_reduce_any(int x) {
if (width == ggml_cuda_get_physical_warp_size()) {
- return __any_sync(0xffffffff, x);
+ return ggml_cuda_any(x);
} else {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x = __shfl_xor_sync(0xffffffff, x, offset, width) || x;
+ x = ggml_cuda_shfl_xor(x, offset, width) || x;
}
return x;
}
@@ -550,7 +579,7 @@ template<int width = WARP_SIZE>
static __device__ __forceinline__ float warp_reduce_max(float x) {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x = fmaxf(x, __shfl_xor_sync(0xffffffff, x, offset, width));
+ x = fmaxf(x, ggml_cuda_shfl_xor(x, offset, width));
}
return x;
}
@@ -560,7 +589,7 @@ static __device__ __forceinline__ T warp_prefix_inclusive_sum(T x) {
const int lane_id = threadIdx.x % width;
#pragma unroll
for (int offset = 1; offset < width; offset <<= 1) {
- const T t = __shfl_up_sync(0xffffffff, x, offset, width);
+ const T t = ggml_cuda_shfl_up(x, offset, width);
if (lane_id >= offset) {
x += t;
}
@@ -573,8 +602,8 @@ static __device__ __forceinline__ float2 warp_prefix_inclusive_sum(float2 a) {
const int lane_id = threadIdx.x % width;
#pragma unroll
for (int offset = 1; offset < width; offset <<= 1) {
- const float t_x = __shfl_up_sync(0xffffffff, a.x, offset, width);
- const float t_y = __shfl_up_sync(0xffffffff, a.y, offset, width);
+ const float t_x = ggml_cuda_shfl_up(a.x, offset, width);
+ const float t_y = ggml_cuda_shfl_up(a.y, offset, width);
if (lane_id >= offset) {
a.x += t_x;
a.y += t_y;
@@ -589,7 +618,7 @@ static __device__ __forceinline__ half2 warp_prefix_inclusive_sum(half2 a) {
const int lane_id = threadIdx.x % width;
#pragma unroll
for (int offset = 1; offset < width; offset <<= 1) {
- const half2 t = __shfl_up_sync(0xffffffff, a, offset, width);
+ const half2 t = ggml_cuda_shfl_up(a, offset, width);
if (lane_id >= offset) {
a = __hadd2(a, t);
}
@@ -716,7 +745,7 @@ template<int width = WARP_SIZE>
static __device__ __forceinline__ half2 warp_reduce_max(half2 x) {
#pragma unroll
for (int offset = width/2; offset > 0; offset >>= 1) {
- x = ggml_cuda_hmax2(x, __shfl_xor_sync(0xffffffff, x, offset, width));
+ x = ggml_cuda_hmax2(x, ggml_cuda_shfl_xor(x, offset, width));
}
return x;
}
@@ -867,7 +896,7 @@ static __device__ __forceinline__ float ggml_cuda_e8m0_to_fp32(uint8_t x) {
}
static __device__ __forceinline__ float ggml_cuda_ue4m3_to_fp32(uint8_t x) {
-#if defined(GGML_USE_HIP) && defined(CDNA3) && defined(FP8_AVAILABLE) && HIP_VERSION >= 60200000
+#if defined(GGML_USE_HIP) && defined(CDNA3) && defined(FP8_AVAILABLE)
// ROCm does not support fp8 in software on devices with fp8 hardware,
// but CDNA3 supports only e4m3_fnuz (no inf).
const uint32_t bits = x * (x != 0x7F && x != 0xFF); // Convert NaN to 0.0f to match CPU implementation.
diff --git a/ggml/src/ggml-cuda/fattn-common.cuh b/ggml/src/ggml-cuda/fattn-common.cuh
index 0e78b9a5b..666542c80 100644
--- a/ggml/src/ggml-cuda/fattn-common.cuh
+++ b/ggml/src/ggml-cuda/fattn-common.cuh
@@ -376,8 +376,8 @@ static __device__ __forceinline__ void quantize_q8_1_to_shared(
}
#pragma unroll
for (int mask = QI8_1/2; mask > 0; mask >>= 1) {
- amax = fmaxf(amax, __shfl_xor_sync(0xFFFFFFFF, amax, mask, 32));
- sum += __shfl_xor_sync(0xFFFFFFFF, sum, mask, 32);
+ amax = fmaxf(amax, ggml_cuda_shfl_xor(amax, mask, 32));
+ sum += ggml_cuda_shfl_xor(sum, mask, 32);
}
const float d = amax / 127;
diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
index f4c096845..2ae59969d 100644
--- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh
+++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
@@ -818,7 +818,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
for (int col = 0; col < cols_per_thread; ++col) {
#pragma unroll
for (int offset = 16; offset >= 4; offset >>= 1) {
- KQ_max_new[col] = fmaxf(KQ_max_new[col], __shfl_xor_sync(0xFFFFFFFF, KQ_max_new[col], offset, warp_size));
+ KQ_max_new[col] = fmaxf(KQ_max_new[col], ggml_cuda_shfl_xor(KQ_max_new[col], offset, warp_size));
}
}
@@ -911,7 +911,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
#endif // defined(TURING_MMA_AVAILABLE)
#pragma unroll
for (int offset = offset_first; offset >= offset_last; offset >>= 1) {
- KQ_max_new[col] = fmaxf(KQ_max_new[col], __shfl_xor_sync(0xFFFFFFFF, KQ_max_new[col], offset, warp_size));
+ KQ_max_new[col] = fmaxf(KQ_max_new[col], ggml_cuda_shfl_xor(KQ_max_new[col], offset, warp_size));
}
}
@@ -1527,7 +1527,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
for (int col = 0; col < cols_per_thread; ++col) {
#pragma unroll
for (int offset = offset_first; offset >= offset_last; offset >>= 1) {
- KQ_rowsum[col] += __shfl_xor_sync(0xFFFFFFFF, KQ_rowsum[col], offset, warp_size);
+ KQ_rowsum[col] += ggml_cuda_shfl_xor(KQ_rowsum[col], offset, warp_size);
}
}
}
@@ -1710,7 +1710,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
#pragma unroll
for (int offset = np*cols_per_warp/2; offset >= cols_per_warp; offset >>= 1) {
if (offset < warp_size) {
- KQ_cmn = fmaxf(KQ_cmn, __shfl_xor_sync(0xFFFFFFFF, KQ_cmn, offset, warp_size));
+ KQ_cmn = fmaxf(KQ_cmn, ggml_cuda_shfl_xor(KQ_cmn, offset, warp_size));
}
}
@@ -1727,7 +1727,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
#pragma unroll
for (int offset = np*cols_per_warp/2; offset >= cols_per_warp; offset >>= 1) {
if (offset < warp_size) {
- KQ_crs += __shfl_xor_sync(0xFFFFFFFF, KQ_crs, offset, warp_size);
+ KQ_crs += ggml_cuda_shfl_xor(KQ_crs, offset, warp_size);
}
}
}
diff --git a/ggml/src/ggml-cuda/fattn-vec.cuh b/ggml/src/ggml-cuda/fattn-vec.cuh
index 5077b1ce1..09a46187d 100644
--- a/ggml/src/ggml-cuda/fattn-vec.cuh
+++ b/ggml/src/ggml-cuda/fattn-vec.cuh
@@ -316,7 +316,7 @@ static __global__ void flash_attn_ext_vec(
for (int j = 0; j < ncols; ++j) {
#pragma unroll
for (int offset = nthreads_KQ; offset < WARP_SIZE; offset <<= 1) {
- KQ_max_new[j] = fmaxf(KQ_max_new[j], __shfl_xor_sync(0xFFFFFFFF, KQ_max_new[j], offset, WARP_SIZE));
+ KQ_max_new[j] = fmaxf(KQ_max_new[j], ggml_cuda_shfl_xor(KQ_max_new[j], offset, WARP_SIZE));
}
const float KQ_max_scale = expf(KQ_max[j] - KQ_max_new[j]);
KQ_max[j] = KQ_max_new[j];
diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu
index 4e4ed086c..fb542d41c 100644
--- a/ggml/src/ggml-cuda/fattn.cu
+++ b/ggml/src/ggml-cuda/fattn.cu
@@ -60,7 +60,7 @@ static __global__ void flash_attn_mask_to_sparse_indices(
selected |= (!oob || q < q1 - q0) && isfinite(__half2float(mask[q*s31 + i]));
}
}
- selected_warp[item] = (lane_mask_t)__ballot_sync(ballot_mask, (int) selected);
+ selected_warp[item] = (lane_mask_t)ggml_cuda_ballot((int) selected, ballot_mask);
warp_count += popc(selected_warp[item]);
}
diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu
index 78768eace..0ee3da6a3 100644
--- a/ggml/src/ggml-cuda/fwht.cu
+++ b/ggml/src/ggml-cuda/fwht.cu
@@ -34,7 +34,7 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons
#pragma unroll
for (int j = 0; j < el_w; j++) {
const float val = reg[j];
- const float val2 = __shfl_xor_sync(0xFFFFFFFF, val, h, warp_size);
+ const float val2 = ggml_cuda_shfl_xor(val, h, warp_size);
reg[j] = (lane & h) == 0 ? val + val2 : val2 - val;
}
@@ -99,7 +99,7 @@ __global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows
#pragma unroll
for (int j = 0; j < NE; j++) {
const float val = reg[j];
- const float val2 = __shfl_xor_sync(0xFFFFFFFF, val, h, warp_size);
+ const float val2 = ggml_cuda_shfl_xor(val, h, warp_size);
reg[j] = (lane & h) == 0 ? val + val2 : val2 - val;
}
}
diff --git a/ggml/src/ggml-cuda/gated-delta-net-mma.cu b/ggml/src/ggml-cuda/gated-delta-net-mma.cu
index 0684adc1a..5fc9ef330 100644
--- a/ggml/src/ggml-cuda/gated-delta-net-mma.cu
+++ b/ggml/src/ggml-cuda/gated-delta-net-mma.cu
@@ -203,7 +203,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m
const float beta = t < valid ? a.beta[goff + (t0 + t) * a.sb2] : 0.f;
#pragma unroll
for (int offset = 1; offset < 32; offset *= 2) {
- const float previous = __shfl_up_sync(0xffffffff, g, offset, 32);
+ const float previous = ggml_cuda_shfl_up(g, offset, 32);
if (lane >= offset) {
g += previous;
}
@@ -213,7 +213,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m
s.prefix[t] = g;
s.beta[t] = beta;
}
- carry = __shfl_sync(0xffffffff, g, 31, 32);
+ carry = ggml_cuda_shfl(g, 31, 32);
}
}
#pragma unroll
@@ -288,7 +288,7 @@ static __global__ __launch_bounds__(256, BLOCKS) void gdn_single(ggml_cuda_gdn_m
x[j] = col == row ? 1.f : (col < row ? -s.lower[s.lower_index(row, col)] : 0.f);
}
for (int pivot = row - 1; pivot > 0; --pivot) {
- const float xp = __shfl_sync(0xffffffff, x[pivot / 32], pivot % 32, 32);
+ const float xp = ggml_cuda_shfl(x[pivot / 32], pivot % 32, 32);
#pragma unroll
for (int j = 0; j < (C + 31) / 32; ++j) {
const int col = lane + 32 * j;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 540c53614..4950b9ea7 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -3635,10 +3635,6 @@ static const struct {
};
static bool ggml_cuda_topk_moe_group_fusion(const struct ggml_cgraph * cgraph, int node_idx, const ggml_tensor * selection, ggml_cuda_topk_moe_args & args) {
-#if defined(GGML_USE_HIP) && HIP_VERSION < 60300000
- GGML_UNUSED_VARS(cgraph, node_idx, selection, args);
- return false;
-#else
const int n_ops = (int) std::size(topk_moe_group_pattern);
if (node_idx + n_ops > cgraph->n_nodes) {
return false;
@@ -3667,7 +3663,6 @@ static bool ggml_cuda_topk_moe_group_fusion(const struct ggml_cgraph * cgraph, i
args.n_groups = n_groups;
args.n_group_used = n_group_used;
return true;
-#endif // defined(GGML_USE_HIP) && HIP_VERSION < 60300000
}
static bool ggml_cuda_topk_moe_fusion(const struct ggml_cgraph * cgraph, int node_idx, ggml_cuda_topk_moe_args & args) {
diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh
index 524db9d85..325cb9306 100644
--- a/ggml/src/ggml-cuda/mma.cuh
+++ b/ggml/src/ggml-cuda/mma.cuh
@@ -52,8 +52,8 @@ static __device__ __forceinline__ int ggml_cuda_movmatrix(const int x) {
const int shift_low = ((src_j + 0) % 2) * 16;
const int shift_high = ((src_j + 1) % 2) * 16;
- const int ret_low = (__shfl_sync(0xFFFFFFFF, x, src_laneid_low, WARP_SIZE) >> shift_low) & 0x0000FFFF;
- const int ret_high = (__shfl_sync(0xFFFFFFFF, x, src_laneid_high, WARP_SIZE) << shift_high) & 0xFFFF0000;
+ const int ret_low = (ggml_cuda_shfl(x, src_laneid_low, WARP_SIZE) >> shift_low) & 0x0000FFFF;
+ const int ret_high = (ggml_cuda_shfl(x, src_laneid_high, WARP_SIZE) << shift_high) & 0xFFFF0000;
return ret_low | ret_high;
}
@@ -736,7 +736,7 @@ namespace ggml_cuda_mma {
int i = threadIdx.x / 16;
tmp[i] = tile_float.x[l];
i ^= 1;
- tmp[i] = __shfl_xor_sync(0xFFFFFFFF, tile_float.x[l], 16, WARP_SIZE);
+ tmp[i] = ggml_cuda_shfl_xor(tile_float.x[l], 16, WARP_SIZE);
ret.x[l] = make_half2(tmp[0], tmp[1]);
}
return ret;
@@ -767,8 +767,8 @@ namespace ggml_cuda_mma {
// On Volta FP16 and FP32 tiles have a different memory layout,
// for the conversion threads with an offset of 2 need to exchange half their values:
- ret.x[l0/2 + (((threadIdx.x % 4) / 2) ^ 1)] = __shfl_xor_sync(
- 0xFFFFFFFF, ret.x[l0/2 + (((threadIdx.x % 4) / 2) ^ 1)], 2, WARP_SIZE);
+ ret.x[l0/2 + (((threadIdx.x % 4) / 2) ^ 1)] = ggml_cuda_shfl_xor(
+ ret.x[l0/2 + (((threadIdx.x % 4) / 2) ^ 1)], 2, WARP_SIZE);
}
return ret;
}
diff --git a/ggml/src/ggml-cuda/mmid.cu b/ggml/src/ggml-cuda/mmid.cu
index 64de4fb27..6b425a377 100644
--- a/ggml/src/ggml-cuda/mmid.cu
+++ b/ggml/src/ggml-cuda/mmid.cu
@@ -88,7 +88,7 @@ static __global__ void mm_ids_helper(
int it_compact_add_lower = 0;
#pragma unroll
for (int offset = neu_padded; offset < warp_size; offset += neu_padded) {
- const int tmp = __shfl_up_sync(0xFFFFFFFF, it_compact_add_self, offset, warp_size);
+ const int tmp = ggml_cuda_shfl_up(it_compact_add_self, offset, warp_size);
if (threadIdx.x >= static_cast<unsigned int>(offset)) {
it_compact_add_lower += tmp;
}
@@ -99,7 +99,7 @@ static __global__ void mm_ids_helper(
}
// The thread with the highest index in the warp always has the sum over the whole warp, use it to increment all threads:
- it_compact += __shfl_sync(0xFFFFFFFF, it_compact_add_lower + it_compact_add_self, warp_size - 1, warp_size);
+ it_compact += ggml_cuda_shfl(it_compact_add_lower + it_compact_add_self, warp_size - 1, warp_size);
}
}
nex_prev = warp_reduce_sum<warp_size>(nex_prev);
diff --git a/ggml/src/ggml-cuda/pool2d.cu b/ggml/src/ggml-cuda/pool2d.cu
index 83a49753c..6cfd05e06 100644
--- a/ggml/src/ggml-cuda/pool2d.cu
+++ b/ggml/src/ggml-cuda/pool2d.cu
@@ -100,7 +100,7 @@ static __global__ void pool2d_nchw_kernel_warp(
#pragma unroll
for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) {
- const To other = __shfl_xor_sync(0xFFFFFFFF, res, offset, WARP_SIZE);
+ const To other = ggml_cuda_shfl_xor(res, offset, WARP_SIZE);
res = (op == GGML_OP_POOL_MAX) ? max(res, other) : res + other;
}
diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu
index 2acf1f6d0..6483da821 100644
--- a/ggml/src/ggml-cuda/quantize.cu
+++ b/ggml/src/ggml-cuda/quantize.cu
@@ -389,7 +389,7 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
float amax = fabsf(xi);
#pragma unroll
for (int mask = 16; mask > 0; mask >>= 1) {
- amax = fmaxf(amax, __shfl_xor_sync(0xFFFFFFFF, amax, mask, WARP_SIZE));
+ amax = fmaxf(amax, ggml_cuda_shfl_xor(amax, mask, WARP_SIZE));
}
const uint8_t e = compute_e8m0_scale(amax);
@@ -399,10 +399,10 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
#if CUDART_VERSION >= 12080
const float scaled_val = xi * inv_s;
- const float val0 = __shfl_sync(0xFFFFFFFF, scaled_val, base, WARP_SIZE);
- const float val1 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 16, WARP_SIZE);
- const float val2 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 1, WARP_SIZE);
- const float val3 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 17, WARP_SIZE);
+ const float val0 = ggml_cuda_shfl(scaled_val, base, WARP_SIZE);
+ const float val1 = ggml_cuda_shfl(scaled_val, base + 16, WARP_SIZE);
+ const float val2 = ggml_cuda_shfl(scaled_val, base + 1, WARP_SIZE);
+ const float val3 = ggml_cuda_shfl(scaled_val, base + 17, WARP_SIZE);
__nv_fp4x4_e2m1 fp4_packed(make_float4(val0, val1, val2, val3));
packed[b] = *(char2 *) &fp4_packed;
@@ -410,10 +410,10 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
// Fallback: manual FP4 conversion using LUT
const uint8_t q_val = ggml_cuda_float_to_fp4_e2m1(xi, inv_s);
- const uint8_t q_lo_0 = __shfl_sync(0xFFFFFFFF, q_val, base, WARP_SIZE);
- const uint8_t q_lo_1 = __shfl_sync(0xFFFFFFFF, q_val, base + 1, WARP_SIZE);
- const uint8_t q_hi_0 = __shfl_sync(0xFFFFFFFF, q_val, base + 16, WARP_SIZE);
- const uint8_t q_hi_1 = __shfl_sync(0xFFFFFFFF, q_val, base + 17, WARP_SIZE);
+ const uint8_t q_lo_0 = ggml_cuda_shfl(q_val, base, WARP_SIZE);
+ const uint8_t q_lo_1 = ggml_cuda_shfl(q_val, base + 1, WARP_SIZE);
+ const uint8_t q_hi_0 = ggml_cuda_shfl(q_val, base + 16, WARP_SIZE);
+ const uint8_t q_hi_1 = ggml_cuda_shfl(q_val, base + 17, WARP_SIZE);
char2 q;
q.x = (q_hi_0 << 4) | q_lo_0;
@@ -497,7 +497,7 @@ static __global__ void quantize_mmq_q8_1(
// Exchange max. abs. value between vals_per_scale/4 threads.
#pragma unroll
for (int offset = vals_per_scale/8; offset > 0; offset >>= 1) {
- amax = fmaxf(amax, __shfl_xor_sync(0xFFFFFFFF, amax, offset, WARP_SIZE));
+ amax = fmaxf(amax, ggml_cuda_shfl_xor(amax, offset, WARP_SIZE));
}
const float d_inv = 127.0f / amax;
@@ -515,7 +515,7 @@ static __global__ void quantize_mmq_q8_1(
// Calculate sums across vals_per_sum/4 threads.
#pragma unroll
for (int offset = vals_per_sum/8; offset > 0; offset >>= 1) {
- sum_q += __shfl_xor_sync(0xFFFFFFFF, sum_q, offset, WARP_SIZE);
+ sum_q += ggml_cuda_shfl_xor(sum_q, offset, WARP_SIZE);
}
sum = d * float(sum_q);
diff --git a/ggml/src/ggml-cuda/top-k.cu b/ggml/src/ggml-cuda/top-k.cu
index 19cb45d1f..227751eea 100644
--- a/ggml/src/ggml-cuda/top-k.cu
+++ b/ggml/src/ggml-cuda/top-k.cu
@@ -184,10 +184,7 @@ static __global__ void top_k_radix_gather(
constexpr int NBINS = 1 << RADIX_BITS;
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr int nwarps = BLOCK_SIZE / warp_size;
- using lane_mask_t = std::conditional_t<warp_size == 64, uint64_t, uint32_t>;
- // HIP requires a 64-bit mask also for a 32-wide warp
- constexpr uint64_t ballot_mask = ~(lane_mask_t) 0;
- const auto popc = [](const lane_mask_t mask) {
+ const auto popc = [](const ggml_cuda_thread_mask_t mask) {
if constexpr (warp_size == 64) {
return __popcll(mask);
} else {
@@ -249,8 +246,8 @@ static __global__ void top_k_radix_gather(
const int64_t col = tile + lane;
const bool is_tie = col < warp_end && top_k_float_to_ordered(row_src[col]) == threshold;
- const lane_mask_t ties = (lane_mask_t) __ballot_sync(ballot_mask, is_tie);
- const int tie_pos = pos + popc(ties & (((lane_mask_t) 1 << lane) - 1));
+ const ggml_cuda_thread_mask_t ties = ggml_cuda_ballot(is_tie);
+ const int tie_pos = pos + popc(ties & (((ggml_cuda_thread_mask_t) 1 << lane) - 1));
if (is_tie && tie_pos < pos_end) {
row_dst[k - rank + tie_pos] = col;
}
diff --git a/ggml/src/ggml-cuda/topk-moe.cu b/ggml/src/ggml-cuda/topk-moe.cu
index e99b2e84a..d6481baaf 100644
--- a/ggml/src/ggml-cuda/topk-moe.cu
+++ b/ggml/src/ggml-cuda/topk-moe.cu
@@ -20,22 +20,24 @@ template <int experts_per_thread, bool use_limit>
__device__ void softmax_warp_inplace(float (&vals)[experts_per_thread], const int limit, const int lane) {
float max_val = -INFINITY;
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = lane + i * WARP_SIZE;
+ const int idx = lane + i * warp_size;
const bool active = !use_limit || (idx < limit);
if (active) {
max_val = max(max_val, vals[i]);
}
}
- max_val = warp_reduce_max(max_val);
+ max_val = warp_reduce_max<warp_size>(max_val);
float sum = 0.f;
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = lane + i * WARP_SIZE;
+ const int idx = lane + i * warp_size;
const bool active = !use_limit || (idx < limit);
if (active) {
const float val = expf(vals[i] - max_val);
@@ -46,13 +48,13 @@ __device__ void softmax_warp_inplace(float (&vals)[experts_per_thread], const in
}
}
- sum = warp_reduce_sum(sum);
+ sum = warp_reduce_sum<warp_size>(sum);
const float inv_sum = 1.0f / sum;
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = lane + i * WARP_SIZE;
+ const int idx = lane + i * warp_size;
const bool active = !use_limit || (idx < limit);
if (active) {
vals[i] *= inv_sum;
@@ -62,9 +64,10 @@ __device__ void softmax_warp_inplace(float (&vals)[experts_per_thread], const in
template <int experts_per_thread, bool use_limit>
__device__ void sigmoid_warp_inplace(float (&vals)[experts_per_thread], const int limit, const int lane) {
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = lane + i * WARP_SIZE;
+ const int idx = lane + i * warp_size;
const bool active = !use_limit || (idx < limit);
vals[i] = active ? 1.f / (1.f + expf(-vals[i])) : -INFINITY;
}
@@ -72,35 +75,37 @@ __device__ void sigmoid_warp_inplace(float (&vals)[experts_per_thread], const in
template <int experts_per_thread, bool use_limit>
__device__ void sqrt_softplus_warp_inplace(float (&vals)[experts_per_thread], const int limit, const int lane) {
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = lane + i * WARP_SIZE;
+ const int idx = lane + i * warp_size;
const bool active = !use_limit || (idx < limit);
vals[i] = active ? sqrtf(vals[i] > 20.0f ? vals[i] : log1pf(expf(vals[i]))) : -INFINITY;
}
}
-#if !defined(GGML_USE_HIP) || HIP_VERSION >= 60300000
// Set -inf for experts outside the top n_group_used groups. Group score is the sum of its top-2 experts (DeepSeek V3).
template <int n_experts, int experts_per_thread>
__device__ void group_mask_warp_inplace(float (&vals)[experts_per_thread], const int n_groups, const int n_group_used, const int lane) {
const int n_per_group = n_experts / n_groups;
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+
float score = -INFINITY; // lane g holds the score of group g
for (int g = 0; g < n_groups; g++) {
float top1 = -INFINITY;
float top2 = -INFINITY;
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- if ((i * WARP_SIZE + lane) / n_per_group == g) {
+ if ((i * warp_size + lane) / n_per_group == g) {
top2 = fmaxf(top2, fminf(top1, vals[i]));
top1 = fmaxf(top1, vals[i]);
}
}
#pragma unroll
- for (int offset = WARP_SIZE / 2; offset > 0; offset >>= 1) {
- const float o1 = __shfl_xor_sync(0xFFFFFFFF, top1, offset, WARP_SIZE);
- const float o2 = __shfl_xor_sync(0xFFFFFFFF, top2, offset, WARP_SIZE);
+ for (int offset = warp_size / 2; offset > 0; offset >>= 1) {
+ const float o1 = ggml_cuda_shfl_xor(top1, offset, warp_size);
+ const float o2 = ggml_cuda_shfl_xor(top2, offset, warp_size);
top2 = fmaxf(fminf(top1, o1), fmaxf(top2, o2));
top1 = fmaxf(top1, o1);
}
@@ -112,20 +117,30 @@ __device__ void group_mask_warp_inplace(float (&vals)[experts_per_thread], const
// rank of each group, lower index wins ties
int rank = 0;
for (int g = 0; g < n_groups; g++) {
- const float other = __shfl_sync(0xFFFFFFFF, score, g, WARP_SIZE);
+ const float other = ggml_cuda_shfl(score, g, warp_size);
rank += other > score || (other == score && g < lane);
}
- const auto keep = __ballot_sync(0xFFFFFFFFULL, lane < n_groups && rank < n_group_used);
+ const auto keep = ggml_cuda_ballot(lane < n_groups && rank < n_group_used);
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int g = (i * WARP_SIZE + lane) / n_per_group;
+ const int g = (i * warp_size + lane) / n_per_group;
if (g >= n_groups || !((keep >> g) & 1)) {
vals[i] = -INFINITY;
}
}
}
-#endif // !defined(GGML_USE_HIP) || HIP_VERSION >= 60300000
+
+static __device__ constexpr int calc_experts_per_thread(int n_experts, int warp_size)
+{
+ if(n_experts <= warp_size) {
+ return 1;
+ }
+ if(n_experts % warp_size == 0) {
+ return n_experts / warp_size;
+ }
+ return n_experts / warp_size + 1;
+}
/*
This kernel does the following:
@@ -137,7 +152,7 @@ __device__ void group_mask_warp_inplace(float (&vals)[experts_per_thread], const
It is intended as fusion of softmax->top-k->get_rows pipeline for MoE models
*/
template <int n_experts, bool has_bias>
-__launch_bounds__(TOPK_MOE_ROWS_PER_BLOCK * WARP_SIZE, 1)
+__launch_bounds__(TOPK_MOE_ROWS_PER_BLOCK * ggml_cuda_get_physical_warp_size(), 1)
__global__ void topk_moe_cuda(const float * logits,
float * weights,
int32_t * ids,
@@ -147,6 +162,8 @@ __global__ void topk_moe_cuda(const float * logits,
const float clamp_val,
const float scale_val,
const topk_moe_config config) {
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+
#if defined(GGML_USE_MUSA)
// MUSA: every warp of a partially filled block must reach the barrier below.
const int row = MIN(blockIdx.x * blockDim.y + threadIdx.y, n_rows - 1);
@@ -161,7 +178,7 @@ __global__ void topk_moe_cuda(const float * logits,
weights += n_expert_used * row;
ids += n_experts * row;
- constexpr int experts_per_thread = (n_experts > WARP_SIZE) ? n_experts / WARP_SIZE : 1;
+ constexpr int experts_per_thread = calc_experts_per_thread(n_experts, warp_size);
float wt[experts_per_thread];
@@ -173,9 +190,9 @@ __global__ void topk_moe_cuda(const float * logits,
ggml_cuda_pdl_sync();
#pragma unroll
- for (int i = 0; i < n_experts; i += WARP_SIZE) {
+ for (int i = 0; i < n_experts; i += warp_size) {
const int expert = i + threadIdx.x;
- wt[i / WARP_SIZE] = (n_experts % WARP_SIZE == 0 || expert < n_experts) ? logits[expert] : -INFINITY;
+ wt[i / warp_size] = (n_experts % warp_size == 0 || expert < n_experts) ? logits[expert] : -INFINITY;
}
// Weights and IDs can alias logits, so wait until every row in the block reads its logits.
@@ -213,17 +230,16 @@ __global__ void topk_moe_cuda(const float * logits,
selection_wt[i] = -INFINITY;
}
#pragma unroll
- for (int i = 0; i < n_experts; i += WARP_SIZE) {
+ for (int i = 0; i < n_experts; i += warp_size) {
const int expert = i + threadIdx.x;
- selection_wt[i / WARP_SIZE] =
- (n_experts % WARP_SIZE == 0 || expert < n_experts) ? wt[i / WARP_SIZE] + bias[expert] : -INFINITY;
- if (__isnanf(selection_wt[i / WARP_SIZE])) {
- selection_wt[i / WARP_SIZE] = -FLT_MAX;
+ selection_wt[i / warp_size] =
+ (n_experts % warp_size == 0 || expert < n_experts) ? wt[i / warp_size] + bias[expert] : -INFINITY;
+ if (__isnanf(selection_wt[i / warp_size])) {
+ selection_wt[i / warp_size] = -FLT_MAX;
}
}
}
-#if !defined(GGML_USE_HIP) || HIP_VERSION >= 60300000
// without bias, masking wt is safe: the matcher ensures the kept groups hold all used experts
if (config.n_groups > 0) {
if constexpr (has_bias) {
@@ -232,7 +248,6 @@ __global__ void topk_moe_cuda(const float * logits,
group_mask_warp_inplace<n_experts>(wt, config.n_groups, config.n_group_used, threadIdx.x);
}
}
-#endif // !defined(GGML_USE_HIP) || HIP_VERSION >= 60300000
//at this point, each thread holds either a portion of the softmax distribution
//or the raw logits. We do the argmax reduce over n_expert_used, each time marking
@@ -257,8 +272,8 @@ __global__ void topk_moe_cuda(const float * logits,
#pragma unroll
for (int i = 1; i < experts_per_thread; i++) {
- const int expert = threadIdx.x + i * WARP_SIZE;
- if ((n_experts % WARP_SIZE == 0 || expert < n_experts) && selection_wt[i] > max_val_s) {
+ const int expert = threadIdx.x + i * warp_size;
+ if ((n_experts % warp_size == 0 || expert < n_experts) && selection_wt[i] > max_val_s) {
max_val = wt[i];
max_val_s = selection_wt[i];
max_expert = expert;
@@ -266,10 +281,10 @@ __global__ void topk_moe_cuda(const float * logits,
}
#pragma unroll
- for (int mask = WARP_SIZE / 2; mask > 0; mask /= 2) {
- const float val = __shfl_xor_sync(0xFFFFFFFF, max_val, mask, WARP_SIZE);
- const float val_s = __shfl_xor_sync(0xFFFFFFFF, max_val_s, mask, WARP_SIZE);
- const int expert = __shfl_xor_sync(0xFFFFFFFF, max_expert, mask, WARP_SIZE);
+ for (int mask = warp_size / 2; mask > 0; mask /= 2) {
+ const float val = ggml_cuda_shfl_xor(max_val, mask, warp_size);
+ const float val_s = ggml_cuda_shfl_xor(max_val_s, mask, warp_size);
+ const int expert = ggml_cuda_shfl_xor(max_expert, mask, warp_size);
if (val_s > max_val_s || (val_s == max_val_s && expert < max_expert)) {
max_val = val;
max_val_s = val_s;
@@ -277,39 +292,39 @@ __global__ void topk_moe_cuda(const float * logits,
}
}
- if ((max_expert & (WARP_SIZE - 1)) == threadIdx.x) {
- selection_wt[max_expert / WARP_SIZE] = -INFINITY;
+ if ((max_expert & (warp_size - 1)) == threadIdx.x) {
+ selection_wt[max_expert / warp_size] = -INFINITY;
}
} else {
#pragma unroll
for (int i = 1; i < experts_per_thread; i++) {
- const int expert = threadIdx.x + i * WARP_SIZE;
- if ((n_experts % WARP_SIZE == 0 || expert < n_experts) && wt[i] > max_val) {
+ const int expert = threadIdx.x + i * warp_size;
+ if ((n_experts % warp_size == 0 || expert < n_experts) && wt[i] > max_val) {
max_val = wt[i];
max_expert = expert;
}
}
#pragma unroll
- for (int mask = WARP_SIZE / 2; mask > 0; mask /= 2) {
- const float val = __shfl_xor_sync(0xFFFFFFFF, max_val, mask, WARP_SIZE);
- const int expert = __shfl_xor_sync(0xFFFFFFFF, max_expert, mask, WARP_SIZE);
+ for (int mask = warp_size / 2; mask > 0; mask /= 2) {
+ const float val = ggml_cuda_shfl_xor(max_val, mask, warp_size);
+ const int expert = ggml_cuda_shfl_xor(max_expert, mask, warp_size);
if (val > max_val || (val == max_val && expert < max_expert)) {
max_val = val;
max_expert = expert;
}
}
- if ((max_expert & (WARP_SIZE - 1)) == threadIdx.x) {
- wt[max_expert / WARP_SIZE] = -INFINITY;
+ if ((max_expert & (warp_size - 1)) == threadIdx.x) {
+ wt[max_expert / warp_size] = -INFINITY;
}
}
- if ((k & (WARP_SIZE - 1)) == threadIdx.x) {
- output_weights[k / WARP_SIZE] = max_val;
+ if ((k & (warp_size - 1)) == threadIdx.x) {
+ output_weights[k / warp_size] = max_val;
}
- if ((max_expert & (WARP_SIZE - 1)) == threadIdx.x) {
+ if ((max_expert & (warp_size - 1)) == threadIdx.x) {
ids[k] = max_expert;
if (config.with_norm) {
wt_sum += max_val;
@@ -318,7 +333,7 @@ __global__ void topk_moe_cuda(const float * logits,
}
if (config.with_norm) {
- wt_sum = warp_reduce_sum(wt_sum);
+ wt_sum = warp_reduce_sum<warp_size>(wt_sum);
wt_sum = max(wt_sum, clamp_val);
const float inv_sum = 1.0f / wt_sum;
@@ -333,7 +348,7 @@ __global__ void topk_moe_cuda(const float * logits,
#pragma unroll
for (int i = 0; i < experts_per_thread; i++) {
- const int idx = i * WARP_SIZE + threadIdx.x;
+ const int idx = i * warp_size + threadIdx.x;
if (idx < n_expert_used) {
weights[idx] = output_weights[i] * scale_val;
}
@@ -355,15 +370,12 @@ static void launch_topk_moe_cuda(ggml_backend_cuda_context & ctx,
GGML_ASSERT(!(config.with_norm && config.delayed_softmax) &&
"delayed softmax is not supported with weight normalization");
const int rows_per_block = TOPK_MOE_ROWS_PER_BLOCK;
+ const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
dim3 grid_dims((n_rows + rows_per_block - 1) / rows_per_block, 1, 1);
- dim3 block_dims(WARP_SIZE, rows_per_block, 1);
+ dim3 block_dims(warp_size, rows_per_block, 1);
cudaStream_t stream = ctx.stream();
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(grid_dims, block_dims, 0, stream);
-#if defined(GGML_USE_HIP) && HIP_VERSION < 60300000
- GGML_ASSERT(config.n_groups == 0 && "expert groups fusion is not supported on this HIP version");
-#endif
-
switch (n_expert) {
case 1:
ggml_cuda_kernel_launch(topk_moe_cuda<1, has_bias>, launch_params,
diff --git a/ggml/src/ggml-cuda/vendors/cuda.h b/ggml/src/ggml-cuda/vendors/cuda.h
index 86c1ed9ce..8253061cb 100644
--- a/ggml/src/ggml-cuda/vendors/cuda.h
+++ b/ggml/src/ggml-cuda/vendors/cuda.h
@@ -6,6 +6,9 @@
#include <cublasLt.h>
#include <cuda_bf16.h>
#include <cuda_fp16.h>
+#include <cstdint>
+
+typedef uint32_t ggml_cuda_thread_mask_t;
#ifdef GGML_USE_NCCL
#include <nccl.h>
diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h
index da159dea1..b350ebe9c 100644
--- a/ggml/src/ggml-cuda/vendors/hip.h
+++ b/ggml/src/ggml-cuda/vendors/hip.h
@@ -1,15 +1,18 @@
#pragma once
-#define HIP_DISABLE_WARP_SYNC_BUILTINS 1
+#define HIP_ENABLE_WARP_SYNC_BUILTINS 1
#include <hip/hip_runtime.h>
#include <hipblas/hipblas.h>
#include <hip/hip_fp16.h>
#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+#include <cstdint>
#ifdef GGML_USE_NCCL
#include <rccl/rccl.h>
#endif // GGML_USE_NCCL
+typedef uint64_t ggml_cuda_thread_mask_t;
#define CUBLAS_GEMM_DEFAULT HIPBLAS_GEMM_DEFAULT
#define CUBLAS_GEMM_DEFAULT_TENSOR_OP HIPBLAS_GEMM_DEFAULT
@@ -29,11 +32,6 @@
#define CU_MEM_LOCATION_TYPE_DEVICE hipMemLocationTypeDevice
#define CU_MEM_ACCESS_FLAGS_PROT_READWRITE hipMemAccessFlagsProtReadWrite
#define CU_CHECK(fn) {hipError_t err = fn; if(err != hipSuccess) { GGML_ABORT("HipVMM Failure: %s\n", hipGetErrorString(err)); }}
-#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 cublasStrsmBatched hipblasStrsmBatched
#define cublasCreate hipblasCreate
#define cublasDestroy hipblasDestroy
@@ -256,16 +254,8 @@
typedef __hip_bfloat16 nv_bfloat16;
typedef __hip_bfloat162 nv_bfloat162;
-
-#if HIP_VERSION >= 60300000
-#include <hip/hip_fp8.h>
typedef __hip_fp8_e4m3 __nv_fp8_e4m3;
#define FP8_AVAILABLE
-#endif // HIP_VERSION >= 60300000
-
-#if HIP_VERSION < 70000000
-#define __ballot_sync(mask, var) __ballot(var)
-#endif //HIP_VERSION < 70000000
typedef int8_t int8x4_t __attribute__((ext_vector_type(4)));
typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4)));
diff --git a/ggml/src/ggml-cuda/vendors/musa.h b/ggml/src/ggml-cuda/vendors/musa.h
index ff13fa052..bfad2a2e6 100644
--- a/ggml/src/ggml-cuda/vendors/musa.h
+++ b/ggml/src/ggml-cuda/vendors/musa.h
@@ -5,11 +5,14 @@
#include <mublas.h>
#include <musa_bf16.h>
#include <musa_fp16.h>
+#include <cstdint>
#ifdef __MUSA_ARCH__
#define __CUDA_ARCH__ 1300 // GGML_CUDA_CC_RUBIN
#endif // __MUSA_ARCH__
+typedef uint32_t ggml_cuda_thread_mask_t;
+
#define CUBLAS_COMPUTE_16F CUDA_R_16F
#define CUBLAS_COMPUTE_32F CUDA_R_32F
#define CUBLAS_COMPUTE_32F_FAST_16F MUBLAS_COMPUTE_32F_FAST_16F
diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt
index 75d4d60e8..35a898e23 100644
--- a/ggml/src/ggml-hip/CMakeLists.txt
+++ b/ggml/src/ggml-hip/CMakeLists.txt
@@ -51,8 +51,8 @@ if (GGML_HIP_RCCL)
find_package(rccl REQUIRED)
endif()
-if (${hip_VERSION} VERSION_LESS 6.1)
- message(FATAL_ERROR "At least ROCM/HIP V6.1 is required")
+if (${hip_VERSION} VERSION_LESS 6.3)
+ message(FATAL_ERROR "At least ROCM/HIP V6.3 is required")
endif()
message(STATUS "HIP and hipBLAS found")