nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Revert CUDA concat change in #2179 (#2200)

This commit is contained in:
Kawrakow 2026-07-28 12:22:55 +03:00 committed by GitHub
parent f0f6ae4bb0
commit b054a8b983
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 4 additions and 64 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

68
ggml/src/ggml-cuda/concat.cu
View File

@ -194,30 +194,6 @@ static __global__ void concat_f32_non_cont(
}
}
static __global__ void k_concat_simple(int64_t n1, int64_t n, const float * __restrict__ src1, const float * __restrict__ src2,
float * __restrict__ dst) {
int64_t i = int64_t(blockIdx.x)*blockDim.x + threadIdx.x;
if (i >= n) {
return;
}
dst[i] = i < n1 ? src1[i] : src2[i - n1];
}
static __global__ void k_concat_dim0(int ne0, int ne00,
size_t nb01, size_t nb02, size_t nb03,
size_t nb11, size_t nb12, size_t nb13,
size_t nb1, size_t nb2, size_t nb3,
const float * __restrict__ src1, const float * __restrict__ src2, float * __restrict__ dst) {
src1 += blockIdx.x * nb01 + blockIdx.y * nb02 + blockIdx.z * nb03;
src2 += blockIdx.x * nb11 + blockIdx.y * nb12 + blockIdx.z * nb13;
dst += blockIdx.x * nb1 + blockIdx.y * nb2 + blockIdx.z * nb3;
for (int i = threadIdx.x; i < ne0; i += blockDim.x) {
dst[i] = i < ne00 ? src1[i] : src2[i - ne00];
}
}
void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
@ -235,44 +211,13 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) &&
(dim == 3 || (dim == 2 && dst->ne[3] == 1) || (dim == 1 && dst->ne[2]*dst->ne[3] == 1))) {
//printf("%s(%s): using cudaMemcpyAsync\n", __func__, dst->name);
constexpr int k_block_size = 512;
int64_t n1 = ggml_nbytes(src0);
int64_t n2 = ggml_nbytes(src1);
if (n1 % sizeof(float) == 0 && n2 % sizeof(float) == 0) {
n1 /= sizeof(float);
n2 /= sizeof(float);
int64_t n = n1 + n2;
int nblocks = (n + k_block_size - 1)/k_block_size;
k_concat_simple<<<nblocks, k_block_size, 0, ctx.stream()>>>(n1, n,
(const float *)src0->data, (const float *)src1->data, (float *)dst->data);
return;
}
//const size_t size0 = ggml_nbytes(src0);
//const size_t size1 = ggml_nbytes(src1);
CUDA_CHECK(cudaMemcpyAsync((char *)dst->data, src0->data, n1, cudaMemcpyDeviceToDevice, stream));
CUDA_CHECK(cudaMemcpyAsync((char *)dst->data + n1, src1->data, n2, cudaMemcpyDeviceToDevice, stream));
const size_t size0 = ggml_nbytes(src0);
const size_t size1 = ggml_nbytes(src1);
CUDA_CHECK(cudaMemcpyAsync((char *)dst->data, src0->data, size0, cudaMemcpyDeviceToDevice, stream));
CUDA_CHECK(cudaMemcpyAsync((char *)dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream));
return;
}
if (dim == 0 && src0->nb[0] == ggml_type_size(src0->type) && src1->nb[0] == ggml_type_size(src1->type) &&
src0->nb[1] % sizeof(float) == 0 && src1->nb[1] % sizeof(float) == 0) {
auto row_size_src0 = ggml_row_size(dst->type, src0->ne[0]);
auto row_size_src1 = ggml_row_size(dst->type, src1->ne[0]);
auto row_size_dst = ggml_row_size(dst->type, dst->ne[0]);
if (row_size_src0 % sizeof(float) == 0 && row_size_src1 % sizeof(float) == 0 && row_size_dst % sizeof(float) == 0) {
auto ne00_eff = row_size_src0/sizeof(float);
auto ne0_eff = row_size_dst /sizeof(float);
dim3 grid(dst->ne[1], dst->ne[2], dst->ne[3]);
k_concat_dim0<<<grid, CUDA_CONCAT_BLOCK_SIZE, 0, ctx.stream()>>>(ne0_eff, ne00_eff,
src0->nb[1]/sizeof(float), src0->nb[2]/sizeof(float), src0->nb[3]/sizeof(float),
src1->nb[1]/sizeof(float), src1->nb[2]/sizeof(float), src1->nb[3]/sizeof(float),
dst->nb[1]/sizeof(float), dst->nb[2]/sizeof(float), dst->nb[3]/sizeof(float),
(const float *)src0->data, (const float *)src1->data, (float *)dst->data);
return;
}
}
if (dim == 0 && src0->nb[0] == ggml_type_size(src0->type) && src1->nb[0] == ggml_type_size(src1->type) &&
src0->nb[1] % sizeof(float) == 0 && src1->nb[1] % sizeof(float) == 0) {
auto bs = ggml_blck_size(dst->type);
@ -280,7 +225,6 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
auto ne00_eff = (src0->ne[0]/bs)*ts/sizeof(float);
auto ne0_eff = (dst->ne[0]/bs)*ts/sizeof(float);
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
//printf("%s(%s): using dim0 contiguous float version with ne3 = %ld\n", __func__, dst->name, dst->ne[3]);
//if (dst->ne[1] >= 65536 || dst->ne[2] >= 65536) {
// fprintf(stderr, "%s: ne1 = %ld, ne2 = %ld exceed max. blocks when computing %s\n", __func__, dst->ne[1], dst->ne[2], dst->name);
// GGML_ABORT("fatal error");
@ -305,7 +249,6 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
//printf("%s(not contiguous): %s(%s) and %s(%s)\n", __func__, src0->name, ggml_type_name(src0->type), src1->name, ggml_type_name(src1->type));
auto ne10_eff = (src1->ne[0]/bs)*ts/sizeof(float);
dim3 grid_dim(dst->ne[1], dst->ne[2], dst->ne[3]);
//printf("%s(%s): using dim0 non-contiguous float version\n", __func__, dst->name);
concat_f32_non_cont<<<grid_dim, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(
(const char *)src0->data,
(const char *)src1->data,
@ -328,7 +271,6 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
GGML_ASSERT(dst->type == GGML_TYPE_F32);
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst) && dim == 2 && dst->ne[3] > 1 && src1->ne[2] == 1) {
//printf("%s(%s): using contiguous dim2 float\n", __func__, dst->name);
float * dst_d = (float *)dst->data;
float * src0_d = (float *)src0->data;
float * src1_d = (float *)src1->data;
@ -337,7 +279,6 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
}
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
printf("%s(%s): using generic contiguous dim2 float\n", __func__, dst->name);
//if (dst->ne[1] >= 65536 || dst->ne[2] >= 65536) {
// fprintf(stderr, "%s: ne1 = %ld, ne2 = %ld exceed max. blocks when computing %s\n", __func__, dst->ne[1], dst->ne[2], dst->name);
// GGML_ABORT("fatal error");
@ -356,7 +297,6 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
dst->ne[0], dst->ne[1], dst->ne[2], dim, stream);
}
} else {
printf("%s(%s): using generic non-contiguous dim2 float\n", __func__, dst->name);
dim3 grid_dim(dst->ne[1], dst->ne[2], dst->ne[3]);
concat_f32_non_cont<<<grid_dim, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(
(const char *)src0->data,