ClaudeBot/llama.cpp
0
0
Fork 1
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-17 20:31:47 +02:00
Code Issues Packages Projects Releases Wiki Activity

ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (#28846)

Browse Source 
* ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA)

* apply logic to NVIDIA as well

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
This commit is contained in:
thelittlefireman
2026-09-14 00:05:10 +02:00
committed by GitHub
co-authored by Johannes Gäßler
parent 7a16a6ce32
commit ad6c66839a
2 changed files with 16 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/common.cuh
+6
View File
@@ -329,6 +329,12 @@ static bool fp16_mma_hardware_available(const int cc) {
(GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
}
// To be used for feature selection of external libraries, e.g. cuBLAS.
static bool fast_bf16_hardware_available(const int cc) {
return (GGML_CUDA_CC_IS_AMD(cc) && (cc >= GGML_CUDA_CC_RDNA3 || GGML_CUDA_CC_IS_CDNA(cc)))
|| (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE);
}
static bool bf16_mma_hardware_available(const int cc) {
return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) ||
GGML_CUDA_CC_IS_CDNA(cc) || cc >= GGML_CUDA_CC_RDNA3 ||
ggml/src/ggml-cuda/ggml-cuda.cu
+10 -2
View File
@@ -1620,11 +1620,19 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
}
static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
const int cc = ggml_cuda_info().devices[ctx.device].cc;
ggml_type compute_type = src0->type;
if (ggml_is_quantized(compute_type)) {
compute_type = fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc)) {
compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
compute_type = GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_BF16 && !fast_bf16_hardware_available(cc)) {
if (GGML_CUDA_CC_IS_AMD(cc) && src1->ne[1] > 32) {
compute_type = GGML_TYPE_F32;
}
if (GGML_CUDA_CC_IS_NVIDIA(cc) && src1->ne[1] > (cc >= GGML_CUDA_CC_VOLTA ? 8 : 128)) {
compute_type = GGML_TYPE_F32;
}
}
if (dst->op_params[0] == GGML_PREC_F32) {
compute_type = GGML_TYPE_F32;