public_git
/
whisper.cpp
mirror of https://github.com/ggml-org/whisper.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

cuda : relax tensor contiguity requirements for quantized concat (llama/25678) 

* 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 <sszymczy@gmail.com>
This commit is contained in:
fairydreaming 2026-07-15 13:36:32 +02:00 committed by Georgi Gerganov
parent e376ba3ad8
commit 8096ba6a0c
No known key found for this signature in database
GPG Key ID: 449E073F9DC10735
3 changed files with 36 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
ggml/src/ggml-cpu/ops.cpp
View File

@ -2081,8 +2081,8 @@ void ggml_compute_forward_concat(
const ggml_tensor * src1 = dst->src[1]; const ggml_tensor * src1 = dst->src[1];
if (ggml_is_quantized(src0->type)) { if (ggml_is_quantized(src0->type)) {
GGML_ASSERT(ggml_is_contiguous(src0)); GGML_ASSERT(ggml_is_contiguous_rows(src0));
GGML_ASSERT(ggml_is_contiguous(src1)); GGML_ASSERT(ggml_is_contiguous_rows(src1));
GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0);
GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0);
} }

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

@ -141,27 +141,25 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE)
template <typename T> template <typename T>
static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim, cudaStream_t stream) { 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 * src0_d = (const T *) src0->data;
const T * src1_d = (const T *) src1->data; const T * src1_d = (const T *) src1->data;
T * dst_d = (T *) dst->data; T * dst_d = (T *) dst->data;
if (dim != 3) { for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) {
for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { concat_cont_cuda(
concat_cont_cuda( src0_d + i3*(src0->nb[3] / sizeof(T)),
src0_d + i3*(src0->nb[3] / sizeof(T)), src1_d + i3*(src1->nb[3] / sizeof(T)),
src1_d + i3*(src1->nb[3] / sizeof(T)), dst_d + i3*( dst->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(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);
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));
} }
} 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 { } else {
GGML_ASSERT(!ggml_is_quantized(src0->type)); 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); GGML_ASSERT(dst->type == src0->type);
if (ggml_is_quantized(src0->type)) { if (ggml_is_quantized(src0->type)) {
GGML_ASSERT(ggml_is_contiguous(src0)); if (dim == 3) {
GGML_ASSERT(ggml_is_contiguous(src1)); 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(src0->ne[0] % ggml_blck_size(src0->type) == 0);
GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->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<uint8_t>(src0, src1, dst, dim, stream); concat_cuda<uint8_t>(src0, src1, dst, dim, stream);
} else { } else {
GGML_ASSERT(ggml_blck_size(src0->type) == 1); GGML_ASSERT(ggml_blck_size(src0->type) == 1);

14
ggml/src/ggml-cuda/ggml-cuda.cu
View File

@ -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 src0_type = op->src[0]->type;
ggml_type src1_type = op->src[1]->type; ggml_type src1_type = op->src[1]->type;
const int32_t dim = op->op_params[0];
return src0_type == src1_type && return src0_type == src1_type &&
src0_type == op->type && src0_type == op->type &&
( (
( (
ggml_is_quantized(src0_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[0]->ne[0] % ggml_blck_size(src0_type) == 0 &&
op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0 op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0
) || ( ) || (