From 8096ba6a0c12c90468e3f708f1e207bbcf118ae3 Mon Sep 17 00:00:00 2001 From: fairydreaming <166155368+fairydreaming@users.noreply.github.com> Date: Wed, 15 Jul 2026 13:36:32 +0200 Subject: [PATCH] cuda : relax tensor contiguity requirements for quantized concat (llama/25678) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : relax tensor contiguity requirements for quantized concat * tests : add test cases for non-contiguous quantized concat * ggml : relax contiguity requirements for quantized concat --------- Co-authored-by: Stanisław Szymczyk --- ggml/src/ggml-cpu/ops.cpp | 4 ++-- ggml/src/ggml-cuda/concat.cu | 41 ++++++++++++++++++--------------- ggml/src/ggml-cuda/ggml-cuda.cu | 14 +++++++++-- 3 files changed, 36 insertions(+), 23 deletions(-) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 58636aa65..8022ab49d 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -2081,8 +2081,8 @@ void ggml_compute_forward_concat( const ggml_tensor * src1 = dst->src[1]; if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous_rows(src0)); + GGML_ASSERT(ggml_is_contiguous_rows(src1)); GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); } diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu index 276ee64e8..6df89013c 100644 --- a/ggml/src/ggml-cuda/concat.cu +++ b/ggml/src/ggml-cuda/concat.cu @@ -141,27 +141,25 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) template static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim, cudaStream_t stream) { - if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + if (dim != 3 && ggml_is_contiguous_to_3(src0) && ggml_is_contiguous_to_3(src1)) { const T * src0_d = (const T *) src0->data; const T * src1_d = (const T *) src1->data; T * dst_d = (T *) dst->data; - if (dim != 3) { - for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { - concat_cont_cuda( - src0_d + i3*(src0->nb[3] / sizeof(T)), - src1_d + i3*(src1->nb[3] / sizeof(T)), - dst_d + i3*( dst->nb[3] / sizeof(T)), - ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], - ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); - } - } else { - 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)); + for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { + concat_cont_cuda( + src0_d + i3*(src0->nb[3] / sizeof(T)), + src1_d + i3*(src1->nb[3] / sizeof(T)), + dst_d + i3*( dst->nb[3] / sizeof(T)), + ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], + ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); } + } else if (dim == 3 && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + 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)); } else { GGML_ASSERT(!ggml_is_quantized(src0->type)); @@ -208,12 +206,17 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { GGML_ASSERT(dst->type == src0->type); if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + if (dim == 3) { + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + } else { + GGML_ASSERT(ggml_is_contiguous_to_3(src0)); + GGML_ASSERT(ggml_is_contiguous_to_3(src1)); + } GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); - // if tensors are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors + // if first 3 dimensions are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors concat_cuda(src0, src1, dst, dim, stream); } else { GGML_ASSERT(ggml_blck_size(src0->type) == 1); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 0878ab9c0..0e185e849 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4816,13 +4816,23 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g { ggml_type src0_type = op->src[0]->type; ggml_type src1_type = op->src[1]->type; + const int32_t dim = op->op_params[0]; return src0_type == src1_type && src0_type == op->type && ( ( ggml_is_quantized(src0_type) && - ggml_is_contiguous(op->src[0]) && - ggml_is_contiguous(op->src[1]) && + ( + ( + dim == 3 && + ggml_is_contiguous(op->src[0]) && + ggml_is_contiguous(op->src[1]) + ) || ( + dim != 3 && + ggml_is_contiguous_to_3(op->src[0]) && + ggml_is_contiguous_to_3(op->src[1]) + ) + ) && op->src[0]->ne[0] % ggml_blck_size(src0_type) == 0 && op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0 ) || (