ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-05-01 20:01:42 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix #261 (#262)

Browse Source 
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
This commit is contained in:
Kawrakow
2025-03-18 07:44:43 +01:00
committed by GitHub
parent bdcae905c4
commit f4ebf13b6a
1 changed files with 33 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1267,7 +1267,7 @@ static void ggml_cuda_op_mul_mat_cublas(
#ifdef GGML_CUDA_IQK_FORCE_BF16
if (ggml_is_quantized(src0->type) && ggml_is_contiguous(src0) && row_diff == src0->ne[1]) {
to_bf16_cuda_t to_bf16_cuda = ggml_get_to_bf16_cuda(src0->type);
GGML_ASSERT(to_bf16_cuda != nullptr);
if (to_bf16_cuda) {
size_t ne = row_diff*ne00;
ggml_cuda_pool_alloc<nv_bfloat16> src0_as_bf16(ctx.pool(id), ne);
to_bf16_cuda(src0_dd_i, src0_as_bf16.get(), row_diff, ne00, stream);
@@ -1302,6 +1302,7 @@ static void ggml_cuda_op_mul_mat_cublas(
to_fp32_cuda(dst_bf16.get(), dst_dd_i, row_diff, src1_ncols, stream);
return;
}
}
#endif
if (compute_capability >= CC_VOLTA && (src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16 || ggml_is_quantized(src0->type)) && ggml_is_contiguous(src0) && row_diff == src0->ne[1] && dst->op_params[0] == GGML_PREC_DEFAULT) {