mirror of
				https://github.com/ggml-org/llama.cpp.git
				synced 2025-11-04 09:32:00 +00:00 
			
		
		
		
	Force FP32 compute in cuBLAS GEMM
This commit is contained in:
		@@ -1245,37 +1245,19 @@ 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();
 | 
			
		||||
 | 
			
		||||
        const float alpha = 1.0f;
 | 
			
		||||
        const float beta = 0.0f;
 | 
			
		||||
 | 
			
		||||
        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, src0_ptr,  CUDA_R_16F, ne00,
 | 
			
		||||
                            src1_ptr,  CUDA_R_16F, ne10,
 | 
			
		||||
                    &beta,   dst_dd_i, CUDA_R_32F, ldc,
 | 
			
		||||
                    CUBLAS_COMPUTE_32F,
 | 
			
		||||
                    CUBLAS_GEMM_DEFAULT_TENSOR_OP));
 | 
			
		||||
 | 
			
		||||
        if (GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) {
 | 
			
		||||
            const float alpha = 1.0f;
 | 
			
		||||
            const float beta = 0.0f;
 | 
			
		||||
            CUBLAS_CHECK(
 | 
			
		||||
                cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
 | 
			
		||||
                        row_diff, src1_ncols, ne10,
 | 
			
		||||
                        &alpha, src0_ptr,  CUDA_R_16F, ne00,
 | 
			
		||||
                                src1_ptr,  CUDA_R_16F, ne10,
 | 
			
		||||
                        &beta,   dst_dd_i, CUDA_R_32F, ldc,
 | 
			
		||||
                        CUBLAS_COMPUTE_32F,
 | 
			
		||||
                        CUBLAS_GEMM_DEFAULT_TENSOR_OP));
 | 
			
		||||
        } else {
 | 
			
		||||
            ggml_cuda_pool_alloc<half> dst_f16(ctx.pool(id), row_diff*src1_ncols);
 | 
			
		||||
 | 
			
		||||
            const half alpha_f16 = 1.0f;
 | 
			
		||||
            const half beta_f16 = 0.0f;
 | 
			
		||||
 | 
			
		||||
            CUBLAS_CHECK(
 | 
			
		||||
                cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N,
 | 
			
		||||
                        row_diff, src1_ncols, ne10,
 | 
			
		||||
                        &alpha_f16, src0_ptr,      CUDA_R_16F, ne00,
 | 
			
		||||
                                    src1_ptr,      CUDA_R_16F, ne10,
 | 
			
		||||
                        &beta_f16,  dst_f16.get(), CUDA_R_16F, ldc,
 | 
			
		||||
                        CUBLAS_COMPUTE_16F,
 | 
			
		||||
                        CUBLAS_GEMM_DEFAULT_TENSOR_OP));
 | 
			
		||||
 | 
			
		||||
            const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16);
 | 
			
		||||
            to_fp32_cuda(dst_f16.get(), dst_dd_i, row_diff*src1_ncols, stream);
 | 
			
		||||
        }
 | 
			
		||||
    } else {
 | 
			
		||||
        ggml_cuda_pool_alloc<float> src0_ddq_as_f32(ctx.pool(id));
 | 
			
		||||
        ggml_cuda_pool_alloc<float> src1_ddq_as_f32(ctx.pool(id));
 | 
			
		||||
 
 | 
			
		||||
		Reference in New Issue
	
	Block a user