mozempk/ggml-org-llama-cpp-mirror
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-10-11 07:20:33 +02:00
Code Issues Packages Projects Releases Wiki Activity

Gate mmvf_f8x2_e4m3_to_bf162 on FP8 availability, i.e. CTK >= 11.8

Browse Source
This commit is contained in:
Oliver Simons
2026-10-09 21:09:51 +02:00
parent 6ee4762f76
commit 8a3884fee8
1 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mmvf.cu
+3 -3
View File
@@ -8,9 +8,9 @@
template <typename T> template <typename T>
using mmvf_y_t = std::conditional_t<std::is_same_v<T, ggml_fp8_e4m3_t>, nv_bfloat16, float>; using mmvf_y_t = std::conditional_t<std::is_same_v<T, ggml_fp8_e4m3_t>, nv_bfloat16, float>;
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) #if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
static __device__ __forceinline__ nv_bfloat162 mmvf_f8x2_e4m3_to_bf162(__nv_fp8x2_storage_t x) { static __device__ __forceinline__ nv_bfloat162 mmvf_f8x2_e4m3_to_bf162(__nv_fp8x2_storage_t x) {
#if defined(FP8_AVAILABLE) && CUDART_VERSION >= 13020 #if CUDART_VERSION >= 13020
return static_cast<nv_bfloat162>(__nv_cvt_fp8x2_to_bf162raw(x, __NV_E4M3)); return static_cast<nv_bfloat162>(__nv_cvt_fp8x2_to_bf162raw(x, __NV_E4M3));
#else #else
return make_bfloat162( return make_bfloat162(
@@ -333,7 +333,7 @@ static __global__ void mul_mat_vec_f(
#endif #endif
} else if constexpr (std::is_same_v<T, ggml_fp8_e4m3_t>) { } else if constexpr (std::is_same_v<T, ggml_fp8_e4m3_t>) {
const nv_bfloat162 * y2 = (const nv_bfloat162 *) y; const nv_bfloat162 * y2 = (const nv_bfloat162 *) y;
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 #if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
const __nv_fp8x2_storage_t * x2 = reinterpret_cast<const __nv_fp8x2_storage_t *>(x); const __nv_fp8x2_storage_t * x2 = reinterpret_cast<const __nv_fp8x2_storage_t *>(x);
const __nv_fp8x2_storage_t * gate_x2 = reinterpret_cast<const __nv_fp8x2_storage_t *>(gate_x); const __nv_fp8x2_storage_t * gate_x2 = reinterpret_cast<const __nv_fp8x2_storage_t *>(gate_x);
nv_bfloat162 sum_bf[ncols_dst] = {}; nv_bfloat162 sum_bf[ncols_dst] = {};