Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-20 01:31:31 +02:00
Code Issues Packages Projects Releases Wiki Activity

use ggml_cuda_syncwarp

Browse Source
This commit is contained in:
Aman Gupta
2026-09-06 16:08:55 +08:00
parent 86e7fc4ffc
commit f0f9902d4c
4 changed files with 23 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/common.cuh
+6
View File
@@ -121,6 +121,12 @@
# 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))
static __device__ __forceinline__ void ggml_cuda_syncwarp() {
#ifndef GGML_USE_HIP
__syncwarp();
#endif // GGML_USE_HIP
}
static __device__ __forceinline__ void ggml_cuda_pdl_sync() {
#if defined(GGML_CUDA_USE_PDL) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_HOPPER
cudaGridDependencySynchronize();
ggml/src/ggml-cuda/fattn-vec.cuh
+1 -3
View File
@@ -317,9 +317,7 @@ static __global__ void flash_attn_ext_vec(
#endif // V_DOT2_F32_F16_AVAILABLE
}
#ifndef GGML_USE_HIP
__syncwarp();
#endif // GGML_USE_HIP
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < WARP_SIZE; k0 += V_cols_per_iter) {
ggml/src/ggml-cuda/mmf.cuh
+15 -15
View File
@@ -143,7 +143,7 @@ static __global__ void mul_mat_f(
if (threadIdx.x == 0) {
slot_map[j] = -1;
}
__syncwarp();
ggml_cuda_syncwarp();
if (col_base + j >= ncols_dst_total) {
continue;
@@ -172,12 +172,12 @@ static __global__ void mul_mat_f(
tile_A A[ntA][warp_size / tile_A::J];
#pragma unroll
for (int itA = 0; itA < ntA; ++itA) {
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int i = 0; i < tile_A::I; ++i) {
tile_xy[i*tile_k_padded + threadIdx.x] = x[(itA*tile_A::I + i)*stride_row + col];
}
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < warp_size; k0 += tile_A::J) {
load_ldmatrix(A[itA][k0/tile_A::J], tile_xy + k0, tile_k_padded);
@@ -186,7 +186,7 @@ static __global__ void mul_mat_f(
#pragma unroll
for (int itB = 0; itB < ntB; ++itB) {
__syncwarp();
ggml_cuda_syncwarp();
if constexpr (std::is_same_v<T, float>) {
#pragma unroll
for (int j0 = 0; j0 < tile_B::I; ++j0) {
@@ -216,7 +216,7 @@ static __global__ void mul_mat_f(
} else {
static_assert(std::is_same_v<T, void>, "unsupported type");
}
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < warp_size; k0 += tile_B::J) {
tile_B B;
@@ -235,7 +235,7 @@ static __global__ void mul_mat_f(
if (nwarps > 1) {
__syncthreads();
} else {
__syncwarp();
ggml_cuda_syncwarp();
}
#pragma unroll
for (int itB = 0; itB < ntB; ++itB) {
@@ -253,7 +253,7 @@ static __global__ void mul_mat_f(
if (nwarps > 1) {
__syncthreads();
} else {
__syncwarp();
ggml_cuda_syncwarp();
}
#pragma unroll
@@ -391,12 +391,12 @@ static __global__ void mul_mat_f_ids(
tile_A A[ntA][warp_size / tile_A::J];
#pragma unroll
for (int itA = 0; itA < ntA; ++itA) {
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int i = 0; i < tile_A::I; ++i) {
tile_xy[i*tile_k_padded + threadIdx.x] = x[(itA*tile_A::I + i)*stride_row + col];
}
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < warp_size; k0 += tile_A::J) {
load_ldmatrix(A[itA][k0/tile_A::J], tile_xy + k0, tile_k_padded);
@@ -430,7 +430,7 @@ static __global__ void mul_mat_f_ids(
int next_buf = 1;
#pragma unroll
for (int itB = 0; itB < ntB; ++itB) {
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int j0 = 0; j0 < tile_B::I; ++j0) {
tile_xy[j0*tile_k_padded + threadIdx.x] = vals_buf[curr_buf][j0];
@@ -440,7 +440,7 @@ static __global__ void mul_mat_f_ids(
gather_tile(itB + 1, vals_buf[next_buf]);
}
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < warp_size; k0 += tile_B::J) {
tile_B B;
@@ -485,7 +485,7 @@ static __global__ void mul_mat_f_ids(
int next_buf = 1;
#pragma unroll
for (int itB = 0; itB < ntB; ++itB) {
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int j0 = 0; j0 < tile_B::I; ++j0) {
const float2 tmp = vals_buf[curr_buf][j0];
@@ -496,7 +496,7 @@ static __global__ void mul_mat_f_ids(
gather_tile(itB + 1, vals_buf[next_buf]);
}
__syncwarp();
ggml_cuda_syncwarp();
#pragma unroll
for (int k0 = 0; k0 < warp_size; k0 += tile_B::J) {
tile_B B;
@@ -523,7 +523,7 @@ static __global__ void mul_mat_f_ids(
if (nwarps > 1) {
__syncthreads();
} else {
__syncwarp();
ggml_cuda_syncwarp();
}
#pragma unroll
for (int itB = 0; itB < ntB; ++itB) {
@@ -541,7 +541,7 @@ static __global__ void mul_mat_f_ids(
if (nwarps > 1) {
__syncthreads();
} else {
__syncwarp();
ggml_cuda_syncwarp();
}
#pragma unroll
ggml/src/ggml-cuda/mmid.cu
+1 -1
View File
@@ -101,7 +101,7 @@ static __global__ void mm_ids_helper(
}
}
nex_prev = warp_reduce_sum<warp_size>(nex_prev);
__syncwarp();
ggml_cuda_syncwarp();
for (int itc = threadIdx.x; itc < it_compact; itc += warp_size) {
const mm_ids_helper_store store_it = store[itc];