From b054a8b983827c01aec59d4dc273a27c492c51c4 Mon Sep 17 00:00:00 2001 From: Kawrakow Date: Tue, 28 Jul 2026 12:22:55 +0300 Subject: [PATCH] Revert CUDA concat change in #2179 (#2200) --- ggml/src/ggml-cuda/concat.cu | 68 +++--------------------------------- 1 file changed, 4 insertions(+), 64 deletions(-) diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu index 7f8069b4..c8695f2c 100644 --- a/ggml/src/ggml-cuda/concat.cu +++ b/ggml/src/ggml-cuda/concat.cu @@ -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<<>>(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<<>>(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<<>>( (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<<>>( (const char *)src0->data,