nikola/llama_cpp
1
0
Fork 0
Code Issues Pull Requests Actions 22 Packages Projects Releases Wiki Activity

cuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels (#19999)

Browse Source
This commit is contained in:
oobabooga
2026-03-01 13:40:22 +08:00
committed by GitHub
parent 05728db18e
commit 66d65ec29b
1 changed files with 28 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/convert.cu
+6 -6
View File
@@ -16,8 +16,7 @@ static __global__ void dequantize_block(const void * __restrict__ vx, dst_t * __
return; return;
} }
const int64_t i01 = blockIdx.y; for (int64_t i01 = blockIdx.y; i01 < ne01; i01 += gridDim.y) {
for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) { for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) {
const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02); const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02);
const int64_t i02 = dm.y; const int64_t i02 = dm.y;
@@ -38,6 +37,7 @@ static __global__ void dequantize_block(const void * __restrict__ vx, dst_t * __
y[iy0 + 0] = ggml_cuda_cast<dst_t>(v.x); y[iy0 + 0] = ggml_cuda_cast<dst_t>(v.x);
y[iy0 + y_offset] = ggml_cuda_cast<dst_t>(v.y); y[iy0 + y_offset] = ggml_cuda_cast<dst_t>(v.y);
} }
}
} }
template <bool need_check> template <bool need_check>
@@ -492,7 +492,7 @@ static void dequantize_block_cuda(const void * vx, dst_t * y,
const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) { const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) {
const int64_t ne0203 = ne02*ne03; const int64_t ne0203 = ne02*ne03;
const uint3 ne02_fdv = init_fastdiv_values(ne02); const uint3 ne02_fdv = init_fastdiv_values(ne02);
const dim3 num_blocks((ne00 + 2*CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / (2*CUDA_DEQUANTIZE_BLOCK_SIZE), ne01, (int)std::min(ne0203, (int64_t)65535)); const dim3 num_blocks((ne00 + 2*CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / (2*CUDA_DEQUANTIZE_BLOCK_SIZE), (int)std::min(ne01, (int64_t)65535), (int)std::min(ne0203, (int64_t)65535));
dequantize_block<qk, qr, dequantize_kernel><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>> dequantize_block<qk, qr, dequantize_kernel><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>
(vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03); (vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03);
} }
@@ -628,10 +628,9 @@ static __global__ void convert_unary(
return; return;
} }
const int64_t i01 = blockIdx.y;
const src_t * x = (const src_t *) vx; const src_t * x = (const src_t *) vx;
for (int64_t i01 = blockIdx.y; i01 < ne01; i01 += gridDim.y) {
for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) { for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) {
const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02); const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02);
const int64_t i02 = dm.y; const int64_t i02 = dm.y;
@@ -641,6 +640,7 @@ static __global__ void convert_unary(
const int64_t iy = (i0203*ne01 + i01)*ne00 + i00; const int64_t iy = (i0203*ne01 + i01)*ne00 + i00;
y[iy] = ggml_cuda_cast<dst_t>(x[ix]); y[iy] = ggml_cuda_cast<dst_t>(x[ix]);
} }
}
} }
template <typename src_t, typename dst_t> template <typename src_t, typename dst_t>
@@ -649,7 +649,7 @@ static void convert_unary_cuda(const void * vx, dst_t * y,
const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) { const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) {
const int64_t ne0203 = ne02*ne03; const int64_t ne0203 = ne02*ne03;
const uint3 ne02_fdv = init_fastdiv_values(ne02); const uint3 ne02_fdv = init_fastdiv_values(ne02);
const dim3 num_blocks((ne00 + CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / CUDA_DEQUANTIZE_BLOCK_SIZE, ne01, (int)std::min(ne0203, (int64_t)65535)); const dim3 num_blocks((ne00 + CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / CUDA_DEQUANTIZE_BLOCK_SIZE, (int)std::min(ne01, (int64_t)65535), (int)std::min(ne0203, (int64_t)65535));
convert_unary<src_t><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>> convert_unary<src_t><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>
(vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03); (vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03);
} }