jdelony/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-06-28 04:30:15 -05:00
Code Issues Packages Projects Releases Wiki Activity

fix: use mmq for volta quantized matmuls (#1785)

Browse Source
This commit is contained in:
Jun Yamog 2026-05-15 17:11:49 +12:00 committed by GitHub
parent 40b65d8f54
commit 53cd4d0ff0
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 6 additions and 29 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1604,33 +1604,6 @@ static void ggml_cuda_op_mul_mat_cublas(
}
const half * src1_ptr = src1->type == GGML_TYPE_F16 ? (const half *) src1_ddf_i : src1_as_f16.get();
// On Volta, avoid storing f32 graph outputs in a temporary f16 buffer;
// finite matmul results outside fp16 range would become +/-inf there.
const bool sm70_f32_output =
compute_capability <= CC_VOLTA &&
dst->type == GGML_TYPE_F32;
if (sm70_f32_output) {
const float alpha_f32 = 1.0f;
const float beta_f32 = 0.0f;
static std::atomic<int> sm70_f32_output_logs{0};
if (sm70_f32_output_logs.fetch_add(1) < 8) {
GGML_CUDA_LOG_WARN(
"%s: using f32 cublas output for %s on cc=%d to avoid fp16 output saturation\n",
__func__, dst->name, compute_capability);
}
CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(id), stream));
CUBLAS_CHECK(
cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
row_diff, src1_ncols, ne10,
&alpha_f32, src0_ptr, CUDA_R_16F, ne00,
src1_ptr, CUDA_R_16F, ne10,
&beta_f32, dst_dd_i, CUDA_R_32F, ldc,
CUBLAS_COMPUTE_32F,
CUBLAS_GEMM_DEFAULT_TENSOR_OP));
return;
}
ggml_cuda_pool_alloc<half> dst_f16(ctx.pool(id), row_diff*src1_ncols);
const half alpha_f16 = 1.0f;

4
ggml/src/ggml-cuda/mmq.cu
View File

@ -247,7 +247,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
#endif //GGML_CUDA_FORCE_MMQ
if (cc < CC_OFFSET_AMD) {
return cc < CC_VOLTA || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
// On Volta, large-batch quantized matmuls otherwise fall back through
// fp16 cuBLAS temporaries. Keep using MMQ for pre-Turing NVIDIA.
return cc < CC_TURING || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
}
return cc < CC_RDNA3 || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;

4
ggml/src/ggml-cuda/mmq_id.cu
View File

@ -566,7 +566,9 @@ bool ggml_cuda_can_use_mmq_id(enum ggml_type type, int cc, int64_t ne11) {
#endif //GGML_CUDA_FORCE_MMQ
if (GGML_CUDA_CC_IS_NVIDIA(cc)) {
return !fp16_mma_hardware_available(cc) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
// Match the plain MMQ policy: use MMQ for pre-Turing NVIDIA, including
// Volta, so indexed/expert matmuls avoid the fp16 cuBLAS fallback.
return cc < CC_TURING || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
}
if (amd_mfma_available(cc)) {