Commit cd4ed902e for llama.cpp
commit cd4ed902e174e54d7198105a720bb379a526309d
Author: Daniel Bevenius <daniel.bevenius@gmail.com>
Date: Tue Nov 10 09:26:53 2026 -0800
cuda : remove unused ids and type_src0 parameters (#32018)
This commit removes the unused ids and type_src0 parameters from
quantize_row_q8_1_cuda.
The motivation for this is that these parameters were required in the
past because this function was used in combination with the function
pointer quantize_cuda_t, which has been now been removed, and this
function no longer needs to adher to the function pointer signature.
Refs: https://github.com/ggml-org/llama.cpp/pull/31884.
diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu
index 8f79123af..b57b26bee 100644
--- a/ggml/src/ggml-cuda/mmvq.cu
+++ b/ggml/src/ggml-cuda/mmvq.cu
@@ -1555,7 +1555,7 @@ void ggml_cuda_mul_mat_vec_q(
const int64_t s11 = src1->nb[1] / ts_src1;
const int64_t s12 = src1->nb[2] / ts_src1;
const int64_t s13 = src1->nb[3] / ts_src1;
- quantize_row_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
+ quantize_row_q8_1_cuda(src1_d, src1_q8_1.get(), ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
}
const int64_t s01 = src0->nb[1] / ts_src0;
diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu
index 727764cdd..4ec8b774b 100644
--- a/ggml/src/ggml-cuda/quantize.cu
+++ b/ggml/src/ggml-cuda/quantize.cu
@@ -557,10 +557,9 @@ static __global__ void quantize_mmq_q8_1(
}
void quantize_row_q8_1_cuda(
- const float * x, const int32_t * ids, void * vy, const ggml_type type_src0,
+ const float * x, void * vy,
const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) {
- GGML_ASSERT(!ids);
GGML_ASSERT(ne0 % QK8_1 == 0);
const uint3 ne2_fastdiv = init_fastdiv_values(ne2);
@@ -570,7 +569,6 @@ void quantize_row_q8_1_cuda(
const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1);
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(num_blocks, block_size, 0, stream);
ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv);
- GGML_UNUSED(type_src0);
}
void quantize_mmq_q8_1_cuda(
diff --git a/ggml/src/ggml-cuda/quantize.cuh b/ggml/src/ggml-cuda/quantize.cuh
index cd761f10a..9ae72ee44 100644
--- a/ggml/src/ggml-cuda/quantize.cuh
+++ b/ggml/src/ggml-cuda/quantize.cuh
@@ -12,8 +12,8 @@ static_assert(MATRIX_ROW_PADDING % CUDA_QUANTIZE_BLOCK_SIZE == 0, "Risk
static_assert(MATRIX_ROW_PADDING % (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ) == 0, "Risk of out-of-bounds access.");
void quantize_row_q8_1_cuda(
- const float * x, const int32_t * ids, void * vy,
- ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03,
+ const float * x, void * vy,
+ int64_t ne00, int64_t s01, int64_t s02, int64_t s03,
int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream);
void quantize_mmq_q8_1_cuda(