Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 01:04:55 +02:00
Code Issues Packages Projects Releases Wiki Activity

Initial CUDA WIP version

Browse Source
This commit is contained in:
Oliver Simons
2026-09-04 21:34:42 +02:00
parent cef6c528f8
commit 38cbaa4ead
31 changed files with 626 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/CMakeLists.txt
+8 -3
View File
@@ -121,7 +121,12 @@ if (CUDAToolkit_FOUND)
template-instances/fattn-vec-instance-f16-f16.cu
template-instances/fattn-vec-instance-q4_0-q4_0.cu
template-instances/fattn-vec-instance-q8_0-q8_0.cu
template-instances/fattn-vec-instance-bf16-bf16.cu)
template-instances/fattn-vec-instance-bf16-bf16.cu
template-instances/fattn-vec-instance-f16-f8_e4m3.cu
template-instances/fattn-vec-instance-f8_e4m3-f16.cu
template-instances/fattn-vec-instance-bf16-f8_e4m3.cu
template-instances/fattn-vec-instance-f8_e4m3-bf16.cu
template-instances/fattn-vec-instance-f8_e4m3-f8_e4m3.cu)
endif()
ggml_add_backend_library(ggml-cuda
@@ -156,7 +161,7 @@ if (CUDAToolkit_FOUND)
if (GGML_STATIC)
if (WIN32)
# As of 12.3.1 CUDA Toolkit for Windows does not offer a static cublas library
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart_static CUDA::cublas)
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart_static CUDA::cublas CUDA::cublasLt)
else ()
if (GGML_CUDA_CUB_3DOT2)
target_link_libraries(ggml-cuda PRIVATE CCCL::CCCL)
@@ -171,7 +176,7 @@ if (CUDAToolkit_FOUND)
if (GGML_CUDA_CUB_3DOT2)
target_link_libraries(ggml-cuda PRIVATE CCCL::CCCL)
endif()
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart CUDA::cublas)
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart CUDA::cublas CUDA::cublasLt)
endif()
if (GGML_CUDA_NO_VMM)
ggml/src/ggml-cuda/common.cuh
+79
View File
@@ -363,6 +363,11 @@ static bool blackwell_mma_available(const int cc) {
ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_RUBIN;
}
static bool fp8_mma_hardware_available(const int cc) {
return GGML_CUDA_CC_IS_NVIDIA(cc) && (cc == GGML_CUDA_CC_ADA_LOVELACE ||
(cc >= GGML_CUDA_CC_BLACKWELL && cc < GGML_CUDA_CC_RUBIN));
}
// Checks whether the tensor's base data pointer and higher-dimensional strides are byte-aligned to `alignment` bytes.
static bool ggml_cuda_is_aligned(const ggml_tensor * tensor, const size_t alignment) {
GGML_ASSERT(tensor != nullptr);
@@ -867,6 +872,60 @@ static __device__ __forceinline__ float ggml_cuda_ue4m3_to_fp32(uint8_t x) {
#endif // defined(GGML_USE_HIP) && defined(CDNA3) && defined(FP8_AVAILABLE) && HIP_VERSION >= 60200000
}
static __device__ __forceinline__ float ggml_cuda_f8_e4m3_to_fp32(uint8_t x) {
const uint8_t magnitude = x & 0x7F;
if (magnitude == 0x7F) {
return NAN;
}
const int exp = (magnitude >> 3) & 0x0F;
const int man = magnitude & 0x07;
float value;
if (exp == 0) {
value = ldexpf((float) man, -9);
} else {
value = ldexpf(1.0f + (float) man / 8.0f, exp - 7);
}
return x & 0x80 ? -value : value;
}
static __device__ __forceinline__ int ggml_cuda_round_to_nearest_even(float x) {
const int value = (int) floorf(x);
const float fraction = x - value;
return fraction > 0.5f || (fraction == 0.5f && (value & 1)) ? value + 1 : value;
}
static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_f8_e4m3(float x) {
const uint8_t sign = signbit(x) ? 0x80 : 0;
x = fabsf(x);
if (isnan(x)) {
return sign | 0x7F;
}
if (x == 0.0f) {
return sign;
}
if (isinf(x) || x >= 448.0f) {
return sign | 0x7E;
}
if (x < 0.015625f) {
return sign | (uint8_t) ggml_cuda_round_to_nearest_even(x * 512.0f);
}
int exp;
const float mantissa = frexpf(x, &exp) * 2.0f;
int encoded_exp = exp + 6;
int encoded_man = ggml_cuda_round_to_nearest_even((mantissa - 1.0f) * 8.0f);
if (encoded_man == 8) {
encoded_man = 0;
encoded_exp++;
}
if (encoded_exp > 15 || (encoded_exp == 15 && encoded_man > 6)) {
return sign | 0x7E;
}
return sign | (uint8_t) (encoded_exp << 3) | (uint8_t) encoded_man;
}
static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) {
#if defined(BLACKWELL_MMA_AVAILABLE) // This is used for NVFP4 subblock scale quantizations only
if (!(x > 0.0f)) {
@@ -1035,6 +1094,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_NVFP4> {
static constexpr int qi = QI_NVFP4;
};
template<>
struct ggml_cuda_type_traits<GGML_TYPE_F8_E4M3> {
static constexpr int qk = QK8_1;
static constexpr int qr = QR8_1;
static constexpr int qi = QI8_1;
};
template<>
struct ggml_cuda_type_traits<GGML_TYPE_Q2_K> {
static constexpr int qk = QK_K;
@@ -1422,6 +1488,9 @@ struct ggml_backend_cuda_context {
cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr};
void * cublas_workspaces[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr};
size_t cublas_workspace_sizes[GGML_CUDA_MAX_DEVICES] = {0};
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
cublasLtHandle_t cublaslt_handles[GGML_CUDA_MAX_DEVICES] = {nullptr};
#endif
int curr_stream_no = 0;
@@ -1516,6 +1585,16 @@ struct ggml_backend_cuda_context {
return cublas_handles[device][curr_stream_no];
}
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
cublasLtHandle_t cublaslt_handle() {
if (cublaslt_handles[device] == nullptr) {
ggml_cuda_set_device(device);
CUBLAS_CHECK(cublasLtCreate(&cublaslt_handles[device]));
}
return cublaslt_handles[device];
}
#endif
// pool
std::unique_ptr<ggml_cuda_pool> pools[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS];
ggml/src/ggml-cuda/convert.cu
+12
View File
@@ -503,6 +503,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cont_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_F32:
return convert_unary_cont_cuda<float>;
case GGML_TYPE_F16:
@@ -563,6 +565,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cont_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_F32:
return convert_unary_cont_cuda<float>;
case GGML_TYPE_BF16:
@@ -620,6 +624,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cont_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_F16:
return convert_unary_cont_cuda<half>;
case GGML_TYPE_BF16:
@@ -647,6 +653,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_BF16:
return convert_unary_cuda<nv_bfloat16>;
default:
@@ -672,6 +680,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_F16:
return convert_unary_cuda<half, nv_bfloat16>;
default:
@@ -697,6 +707,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return dequantize_block_cuda<1, 1, dequantize_f8_e4m3>;
case GGML_TYPE_BF16:
return convert_unary_cuda<nv_bfloat16, float>;
default:
ggml/src/ggml-cuda/dequantize.cuh
+8
View File
@@ -119,6 +119,14 @@ static __device__ __forceinline__ void dequantize_q8_0(const void * vx, const in
v.y *= d;
}
static __device__ __forceinline__ void dequantize_f8_e4m3(const void * vx, const int64_t ib, const int iqs, float2 & v) {
const ggml_fp8_e4m3_t * x = (const ggml_fp8_e4m3_t *) vx;
v.x = ggml_cuda_f8_e4m3_to_fp32(x[ib + 0].bits);
v.y = ggml_cuda_f8_e4m3_to_fp32(x[ib + 1].bits);
GGML_UNUSED(iqs);
}
//================================== k-quants
// Each call dequantizes one super-block of QK_K values into y using the
ggml/src/ggml-cuda/fattn-common.cuh
+51
View File
@@ -145,6 +145,34 @@ static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_bf16(
return sum;
}
template <int D, int nthreads>
static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_f8_e4m3(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
const uint8_t * K_f8 = (const uint8_t *) K_c;
GGML_UNUSED(Q_q8);
GGML_UNUSED(Q_ds_v);
float sum = 0.0f;
#pragma unroll
for (int k_KQ_0 = 0; k_KQ_0 < D; k_KQ_0 += nthreads*4) {
uint8_t tmp[4];
ggml_cuda_memcpy_1<sizeof(tmp)>(tmp, K_f8 + k_KQ_0 + (threadIdx.x % nthreads)*4);
#pragma unroll
for (int k_KQ_1 = 0; k_KQ_1 < 4; ++k_KQ_1) {
const float k = ggml_cuda_f8_e4m3_to_fp32(tmp[k_KQ_1]);
#ifdef V_DOT2_F32_F16_AVAILABLE
sum += k * __half2float(((const half *) Q_v)[k_KQ_0/nthreads + k_KQ_1]);
#else
sum += k * ((const float *) Q_v)[k_KQ_0/nthreads + k_KQ_1];
#endif // V_DOT2_F32_F16_AVAILABLE
}
}
return sum;
}
template<int D, int nthreads>
static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_q4_0(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
@@ -405,6 +433,25 @@ static __device__ __forceinline__ void dequantize_V_bf16(const void * __restrict
}
}
template <typename T, int ne>
static __device__ __forceinline__ void dequantize_V_f8_e4m3(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
static_assert(ne == 2 || ne == 4, "bad ne");
uint8_t tmp[ne];
ggml_cuda_memcpy_1<ne>(tmp, (const uint8_t *) vx + i0);
#pragma unroll
for (int l = 0; l < ne; ++l) {
const float value = ggml_cuda_f8_e4m3_to_fp32(tmp[l]);
if constexpr (std::is_same_v<T, half>) {
((half *) dst)[l] = __float2half(value);
} else if constexpr (std::is_same_v<T, float>) {
((float *) dst)[l] = value;
} else {
static_assert(std::is_same_v<T, void>, "unsupported type");
}
}
}
template <typename T, int ne>
static __device__ __forceinline__ void dequantize_V_q4_0(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
const block_q4_0 * x = (const block_q4_0 *) vx;
@@ -633,6 +680,8 @@ constexpr __device__ vec_dot_KQ_t get_vec_dot_KQ() {
return vec_dot_fattn_vec_KQ_q8_0<D, nthreads>;
} else if constexpr (type_K == GGML_TYPE_BF16) {
return vec_dot_fattn_vec_KQ_bf16<D, nthreads>;
} else if constexpr (type_K == GGML_TYPE_F8_E4M3) {
return vec_dot_fattn_vec_KQ_f8_e4m3<D, nthreads>;
} else {
static_assert(type_K == -1, "bad type");
return nullptr;
@@ -655,6 +704,8 @@ constexpr __device__ dequantize_V_t get_dequantize_V() {
return dequantize_V_q8_0<T, ne>;
} else if constexpr (type_V == GGML_TYPE_BF16) {
return dequantize_V_bf16<float, ne>;
} else if constexpr (type_V == GGML_TYPE_F8_E4M3) {
return dequantize_V_f8_e4m3<T, ne>;
} else {
static_assert(type_V == -1, "bad type");
return nullptr;
ggml/src/ggml-cuda/fattn-vec.cuh
+4
View File
@@ -585,6 +585,7 @@ void ggml_cuda_flash_attn_ext_vec_case(ggml_backend_cuda_context & ctx, ggml_ten
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q5_1); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q8_0); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_BF16); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_F8_E4M3); \
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q4_0)
@@ -593,6 +594,7 @@ EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_F8_E4M3)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q4_0)
@@ -601,6 +603,7 @@ EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_F8_E4M3)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q4_0)
@@ -609,3 +612,4 @@ EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_F8_E4M3)
ggml/src/ggml-cuda/fattn.cu
+26 -1
View File
@@ -402,6 +402,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
@@ -410,6 +411,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q4_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_1)
@@ -418,6 +420,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q4_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q5_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_0)
@@ -426,6 +429,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q5_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q5_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_1)
@@ -434,6 +438,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q5_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_1)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q8_0)
@@ -442,6 +447,7 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_BF16)
@@ -450,11 +456,26 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_F8_E4M3)
#else
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_F16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_F8_E4M3)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_BF16)
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F8_E4M3, GGML_TYPE_F8_E4M3)
#endif // GGML_CUDA_FA_ALL_QUANTS
GGML_ABORT("fatal error");
@@ -482,6 +503,7 @@ static bool ggml_cuda_fattn_kv_type_supported(ggml_type type) {
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_BF16:
case GGML_TYPE_F8_E4M3:
return true;
default:
return false;
@@ -573,7 +595,10 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
}
#ifndef GGML_CUDA_FA_ALL_QUANTS
if (K->type != V->type) {
const bool mixed_fp8 =
(K->type == GGML_TYPE_F8_E4M3 && (V->type == GGML_TYPE_F32 || V->type == GGML_TYPE_F16 || V->type == GGML_TYPE_BF16)) ||
(V->type == GGML_TYPE_F8_E4M3 && (K->type == GGML_TYPE_F32 || K->type == GGML_TYPE_F16 || K->type == GGML_TYPE_BF16));
if (K->type != V->type && !mixed_fp8) {
return BEST_FATTN_KERNEL_NONE;
}
#endif // GGML_CUDA_FA_ALL_QUANTS
ggml/src/ggml-cuda/fp8.cu
+217
View File
@@ -0,0 +1,217 @@
#include "fp8.cuh"
static __global__ void mul_mat_fp8_fallback(
const char * src0, const char * src1, char * dst, int64_t ne00, int64_t ne01, int64_t ne11,
int64_t ne12, int64_t ne13, int64_t r2, int64_t r3, int64_t nb01, int64_t nb02, int64_t nb03,
int64_t nb11, int64_t nb12, int64_t nb13, int64_t nb1, int64_t nb2, int64_t nb3, int64_t ne_dst) {
for (int64_t id = blockIdx.x; id < ne_dst; id += gridDim.x) {
int64_t tmp = id / ne01;
const int64_t i0 = id - tmp*ne01;
const int64_t i1 = tmp % ne11;
tmp /= ne11;
const int64_t i2 = tmp % ne12;
const int64_t i3 = tmp / ne12;
const ggml_fp8_e4m3_t * x = (const ggml_fp8_e4m3_t *) (src0 + i0*nb01 + (i2/r2)*nb02 + (i3/r3)*nb03);
const float * y = (const float *) (src1 + i1*nb11 + i2*nb12 + i3*nb13);
float sum = 0.0f;
for (int64_t k = threadIdx.x; k < ne00; k += blockDim.x) {
sum = fmaf(ggml_cuda_f8_e4m3_to_fp32(x[k].bits), y[k], sum);
}
__shared__ float shared[WARP_SIZE];
sum = block_reduce<block_reduce_method::SUM, 256>(sum, shared);
if (threadIdx.x == 0) {
*(float *) (dst + i0*sizeof(float) + i1*nb1 + i2*nb2 + i3*nb3) = sum;
}
}
}
void ggml_cuda_mul_mat_fp8_fallback(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0->type == GGML_TYPE_F8_E4M3);
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_ASSERT(src0->nb[0] == sizeof(ggml_fp8_e4m3_t));
GGML_ASSERT(src1->nb[0] == sizeof(float));
const int64_t r2 = src1->ne[2] / src0->ne[2];
const int64_t r3 = src1->ne[3] / src0->ne[3];
const int64_t ne_dst = ggml_nelements(dst);
const int blocks = std::min<int64_t>(ne_dst, 65535);
mul_mat_fp8_fallback<<<blocks, 256, 0, ctx.stream()>>>(
(const char *) src0->data, (const char *) src1->data, (char *) dst->data,
src0->ne[0], src0->ne[1], src1->ne[1], src1->ne[2], src1->ne[3], r2, r3,
src0->nb[1], src0->nb[2], src0->nb[3], src1->nb[1], src1->nb[2], src1->nb[3],
dst->nb[1], dst->nb[2], dst->nb[3], ne_dst);
CUDA_CHECK(cudaGetLastError());
}
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
struct fp8_abs_src {
const float * x;
int64_t ne0;
int64_t ne1;
int64_t ne2;
int64_t s1;
int64_t s2;
int64_t s3;
__device__ float operator()(int64_t i) const {
const int64_t i0 = i % ne0;
i /= ne0;
const int64_t i1 = i % ne1;
i /= ne1;
const int64_t i2 = i % ne2;
const int64_t i3 = i / ne2;
const float value = fabsf(x[i0 + i1*s1 + i2*s2 + i3*s3]);
return isfinite(value) ? value : 448.0f;
}
};
static __global__ void fp8_amax_partials(fp8_abs_src src, int64_t ne, float * partials) {
float amax = 0.0f;
for (int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; i < ne; i += (int64_t) blockDim.x*gridDim.x) {
amax = fmaxf(amax, src(i));
}
__shared__ float shared[WARP_SIZE];
amax = block_reduce<block_reduce_method::MAX, 256>(amax, shared);
if (threadIdx.x == 0) {
partials[blockIdx.x] = amax;
}
}
static __global__ void fp8_amax_final(const float * partials, int n, float * amax) {
float value = 0.0f;
for (int i = threadIdx.x; i < n; i += blockDim.x) {
value = fmaxf(value, partials[i]);
}
__shared__ float shared[WARP_SIZE];
value = block_reduce<block_reduce_method::MAX, 256>(value, shared);
if (threadIdx.x == 0) {
*amax = value;
}
}
static __global__ void quantize_fp8_e4m3(
const float * __restrict__ x, uint8_t * __restrict__ y, const float * __restrict__ amax,
float * __restrict__ scale, int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne, int64_t s1, int64_t s2, int64_t s3) {
const int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x;
if (i >= ne) {
return;
}
int64_t tmp = i / ne0;
const int64_t i0 = i - tmp*ne0;
const int64_t i1 = tmp % ne1;
tmp /= ne1;
const int64_t i2 = tmp % ne2;
const int64_t i3 = tmp / ne2;
const float d = *amax > 0.0f ? *amax / 448.0f : 1.0f;
const __nv_fp8_e4m3 q(x[i0 + i1*s1 + i2*s2 + i3*s3] / d);
y[i] = q.__x;
if (i == 0) {
*scale = d;
}
}
static void fp8_destroy_matmul(
cublasLtMatmulDesc_t op_desc, cublasLtMatrixLayout_t a_desc, cublasLtMatrixLayout_t b_desc,
cublasLtMatrixLayout_t d_desc) {
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(d_desc));
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(b_desc));
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(a_desc));
CUBLAS_CHECK(cublasLtMatmulDescDestroy(op_desc));
}
bool ggml_cuda_mul_mat_fp8(
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;
if (!fp8_mma_hardware_available(cc) || src0->type != GGML_TYPE_F8_E4M3 || src1->type != GGML_TYPE_F32 ||
dst->type != GGML_TYPE_F32 || !ggml_is_contiguous(dst) || src0->ne[0] % 16 != 0 || src0->ne[1] % 16 != 0 ||
src0->nb[0] != sizeof(uint8_t) || src0->nb[1] != (size_t) src0->ne[0] || src1->nb[0] != sizeof(float)) {
return false;
}
GGML_TENSOR_BINARY_OP_LOCALS
GGML_ASSERT(ne10 == ne00);
GGML_ASSERT(ne0 == ne01);
GGML_ASSERT(ne12 % ne02 == 0);
GGML_ASSERT(ne13 % ne03 == 0);
cudaStream_t stream = ctx.stream();
const int64_t ne_src1 = ggml_nelements(src1);
ggml_cuda_pool_alloc<uint8_t> src1_fp8(ctx.pool(), ne_src1);
ggml_cuda_pool_alloc<float> src1_scale(ctx.pool(), 1);
ggml_cuda_pool_alloc<float> src1_amax(ctx.pool(), 1);
const fp8_abs_src abs_src = {
(const float *) src1->data, ne10, ne11, ne12,
(int64_t) (nb11 / sizeof(float)), (int64_t) (nb12 / sizeof(float)), (int64_t) (nb13 / sizeof(float))
};
const int reduce_blocks = std::min<int64_t>((ne_src1 + 255)/256, 1024);
ggml_cuda_pool_alloc<float> reduce_tmp(ctx.pool(), reduce_blocks);
fp8_amax_partials<<<reduce_blocks, 256, 0, stream>>>(abs_src, ne_src1, reduce_tmp.ptr);
fp8_amax_final<<<1, 256, 0, stream>>>(reduce_tmp.ptr, reduce_blocks, src1_amax.ptr);
quantize_fp8_e4m3<<<(ne_src1 + 255)/256, 256, 0, stream>>>(
(const float *) src1->data, src1_fp8.ptr, src1_amax.ptr, src1_scale.ptr, ne10, ne11, ne12, ne_src1,
nb11 / sizeof(float), nb12 / sizeof(float), nb13 / sizeof(float));
CUDA_CHECK(cudaGetLastError());
cublasLtMatmulDesc_t op_desc;
cublasLtMatrixLayout_t a_desc;
cublasLtMatrixLayout_t b_desc;
cublasLtMatrixLayout_t d_desc;
CUBLAS_CHECK(cublasLtMatmulDescCreate(&op_desc, CUBLAS_COMPUTE_32F, CUDA_R_32F));
const cublasOperation_t trans_a = CUBLAS_OP_T;
CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_TRANSA, &trans_a, sizeof(trans_a)));
CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(
op_desc, CUBLASLT_MATMUL_DESC_B_SCALE_POINTER, &src1_scale.ptr, sizeof(src1_scale.ptr)));
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&a_desc, CUDA_R_8F_E4M3, ne00, ne01, ne00));
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&b_desc, CUDA_R_8F_E4M3, ne10, ne11, ne10));
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&d_desc, CUDA_R_32F, ne0, ne1, ne0));
cublasLtMatmulPreference_t preference;
CUBLAS_CHECK(cublasLtMatmulPreferenceCreate(&preference));
cublasLtMatmulHeuristicResult_t heuristic;
int returned = 0;
const cublasStatus_t status = cublasLtMatmulAlgoGetHeuristic(
ctx.cublaslt_handle(), op_desc, a_desc, b_desc, d_desc, d_desc, preference, 1, &heuristic, &returned);
CUBLAS_CHECK(cublasLtMatmulPreferenceDestroy(preference));
if (status != CUBLAS_STATUS_SUCCESS || returned == 0) {
fp8_destroy_matmul(op_desc, a_desc, b_desc, d_desc);
return false;
}
const float alpha = 1.0f;
const float beta = 0.0f;
const int64_t r2 = ne12 / ne02;
const int64_t r3 = ne13 / ne03;
for (int64_t i3 = 0; i3 < ne13; ++i3) {
for (int64_t i2 = 0; i2 < ne12; ++i2) {
const char * a = (const char *) src0->data + (i2/r2)*nb02 + (i3/r3)*nb03;
const uint8_t * b = src1_fp8.ptr + (i3*ne12 + i2)*ne11*ne10;
float * d = (float *) ((char *) dst->data + i2*dst->nb[2] + i3*dst->nb[3]);
CUBLAS_CHECK(cublasLtMatmul(ctx.cublaslt_handle(), op_desc, &alpha, a, a_desc, b, b_desc,
&beta, d, d_desc, d, d_desc, &heuristic.algo, nullptr, 0, stream));
}
}
fp8_destroy_matmul(op_desc, a_desc, b_desc, d_desc);
return true;
}
#else
bool ggml_cuda_mul_mat_fp8(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_UNUSED_VARS(ctx, src0, src1, dst);
return false;
}
#endif
ggml/src/ggml-cuda/fp8.cuh
+9
View File
@@ -0,0 +1,9 @@
#pragma once
#include "common.cuh"
bool ggml_cuda_mul_mat_fp8(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
void ggml_cuda_mul_mat_fp8_fallback(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
ggml/src/ggml-cuda/getrows.cu
+4
View File
@@ -344,6 +344,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
get_rows_cuda_q<QK8_0, QR8_0, dequantize_q8_0>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_F8_E4M3:
get_rows_cuda_q<1, 1, dequantize_f8_e4m3>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_Q2_K:
get_rows_cuda_kq<64, dst_t, dequantize_q2_K<dst_t>>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
ggml/src/ggml-cuda/ggml-cuda.cu
+36 -12
View File
@@ -25,6 +25,7 @@
#include "ggml-cuda/diagmask.cuh"
#include "ggml-cuda/diag.cuh"
#include "ggml-cuda/fattn.cuh"
#include "ggml-cuda/fp8.cuh"
#include "ggml-cuda/fwht.cuh"
#include "ggml-cuda/getrows.cuh"
#include "ggml-cuda/im2col.cuh"
@@ -720,6 +721,11 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() {
CUDA_CHECK(cudaFree(cublas_workspaces[i][j]));
}
}
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
if (cublaslt_handles[i] != nullptr) {
CUBLAS_CHECK(cublasLtDestroy(cublaslt_handles[i]));
}
#endif
}
}
@@ -1795,11 +1801,12 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) {
ggml_nbytes(src0) != ggml_backend_buffer_get_alloc_size(src0->buffer, src0) &&
src0->view_src;
bool use_mul_mat_vec_q = ggml_is_quantized(src0->type) && !bad_padding_clear && src1->type == GGML_TYPE_F32 &&
dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE;
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
bool use_mul_mat_vec_q = ggml_cuda_should_use_mmvq(src0->type, cc, src1->ne[1]) && !bad_padding_clear &&
src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32;
use_mul_mat_vec_q = use_mul_mat_vec_q && (src0->type != GGML_TYPE_F8_E4M3 || src0->ne[0] % QK8_1 == 0);
// fusion is not universally faster on Pascal
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
if (cc <= GGML_CUDA_CC_PASCAL) {
return false;
}
@@ -1860,7 +1867,8 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst);
return;
}
if (ggml_cuda_should_use_mmvq(src0->type, cc, ne11)) {
if (ggml_cuda_should_use_mmvq(src0->type, cc, ne11) &&
(src0->type != GGML_TYPE_F8_E4M3 || ne00 % QK8_1 == 0)) {
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst);
return;
}
@@ -1868,6 +1876,12 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
ggml_cuda_mul_mat_q(ctx, src0, src1, nullptr, dst);
return;
}
if (src0->type == GGML_TYPE_F8_E4M3) {
if (!ggml_cuda_mul_mat_fp8(ctx, src0, src1, dst)) {
ggml_cuda_mul_mat_fp8_fallback(ctx, src0, src1, dst);
}
return;
}
ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
}
@@ -1882,8 +1896,9 @@ static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int c
}
if (dst->ne[2] <= MMVQ_MAX_BATCH_SIZE) {
if (ggml_is_quantized(src0->type)) {
if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc)) {
if (ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3) {
if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc) &&
(src0->type != GGML_TYPE_F8_E4M3 || src0->ne[0] % QK8_1 == 0)) {
return false;
}
} else if (GGML_CUDA_CC_IS_AMD(cc)) {
@@ -1918,9 +1933,10 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
if (src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
static_assert(MMVQ_MAX_BATCH_SIZE == MMVF_MAX_BATCH_SIZE);
if (ne2 <= MMVQ_MAX_BATCH_SIZE) {
if (ggml_is_quantized(src0->type)) {
if (ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3) {
const int mmvq_mmid_max = get_mmvq_mmid_max_batch(src0->type, cc);
if (ne2 <= mmvq_mmid_max) {
if (ne2 <= mmvq_mmid_max &&
(src0->type != GGML_TYPE_F8_E4M3 || ne00 % QK8_1 == 0)) {
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, ids, dst);
return;
}
@@ -1951,7 +1967,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
GGML_ASSERT(nb2 % nb1 == 0);
const ggml_type type_src1_sorted = (src0->type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc))
|| ggml_is_quantized(src0->type) ? GGML_TYPE_F32 : src0->type;
|| ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3 ? GGML_TYPE_F32 : src0->type;
const ggml_type type_dst_sorted = GGML_TYPE_F32;
const size_t ts_src1_sorted = ggml_type_size(type_src1_sorted);
const size_t ts_dst_sorted = ggml_type_size(type_dst_sorted);
@@ -3638,7 +3654,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
}
const ggml_tensor * scale = scale_lhs_mm ? scale_node->src[1] : scale_node->src[0];
if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1 ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
@@ -3658,7 +3674,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
}
const ggml_tensor * scale = reshape->src[0];
if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm_node->src[0]->ne[2] ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
@@ -5151,6 +5167,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_Q8_0:
case GGML_TYPE_MXFP4:
case GGML_TYPE_NVFP4:
case GGML_TYPE_F8_E4M3:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
@@ -5188,6 +5205,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_F8_E4M3:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
@@ -5219,7 +5237,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
{
return (
(
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 || op->type == GGML_TYPE_F8_E4M3 ||
op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q5_0 ||
op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_IQ4_NL) &&
op->src[0]->type == GGML_TYPE_F32
@@ -5649,6 +5667,12 @@ static ggml_backend_feature * ggml_backend_cuda_get_features(ggml_backend_reg_t
{
const auto & info = ggml_cuda_info();
for (int id = 0; id < info.device_count; ++id) {
if (fp8_mma_hardware_available(info.devices[id].cc)) {
features.push_back({ "NATIVE_FP8", "1"});
break;
}
}
for (int id = 0; id < info.device_count; ++id) {
if (blackwell_mma_available(info.devices[id].cc)) {
features.push_back({ "BLACKWELL_NATIVE_FP4", "1"});
ggml/src/ggml-cuda/mmvq.cu
+31 -17
View File
@@ -8,6 +8,10 @@
typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
static constexpr __host__ __device__ bool is_scaled_low_precision_type(ggml_type type) {
return type == GGML_TYPE_NVFP4 || type == GGML_TYPE_F8_E4M3;
}
static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1;
@@ -19,6 +23,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
case GGML_TYPE_Q8_0: return vec_dot_q8_0_q8_1;
case GGML_TYPE_MXFP4: return vec_dot_mxfp4_q8_1;
case GGML_TYPE_NVFP4: return vec_dot_nvfp4_q8_1;
case GGML_TYPE_F8_E4M3: return vec_dot_f8_e4m3_q8_1;
case GGML_TYPE_Q2_K: return vec_dot_q2_K_q8_1;
case GGML_TYPE_Q3_K: return vec_dot_q3_K_q8_1;
case GGML_TYPE_Q4_K: return vec_dot_q4_K_q8_1;
@@ -48,6 +53,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {
case GGML_TYPE_Q8_0: return VDR_Q8_0_Q8_1_MMVQ;
case GGML_TYPE_MXFP4: return VDR_MXFP4_Q8_1_MMVQ;
case GGML_TYPE_NVFP4: return VDR_NVFP4_Q8_1_MMVQ;
case GGML_TYPE_F8_E4M3: return VDR_F8_E4M3_Q8_1_MMVQ;
case GGML_TYPE_Q2_K: return VDR_Q2_K_Q8_1_MMVQ;
case GGML_TYPE_Q3_K: return VDR_Q3_K_Q8_1_MMVQ;
case GGML_TYPE_Q4_K: return VDR_Q4_K_Q8_1_MMVQ;
@@ -287,7 +293,7 @@ int get_mmvq_mmid_max_batch(ggml_type type, int cc) {
}
bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) {
if (!ggml_is_quantized(type)) {
if (!ggml_is_quantized(type) && type != GGML_TYPE_F8_E4M3) {
return false;
}
// k-quants cost more to decode and mvq redoes that per column, so MMQ wins sooner.
@@ -570,6 +576,7 @@ static __global__ void mul_mat_vec_q(
constexpr int qk = ggml_cuda_type_traits<type>::qk;
constexpr int qi = ggml_cuda_type_traits<type>::qi;
constexpr int vdr = get_vdr_mmvq(type);
constexpr int kbx_stride = type == GGML_TYPE_F8_E4M3 ? QK8_1 : 1;
constexpr mmvq_parameter_table_id table_id = get_device_table_id();
constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
@@ -618,7 +625,7 @@ static __global__ void mul_mat_vec_q(
gate_bias = (const float *) fusion.gate_bias;
active_glu = fusion.glu_op;
glu_limit = fusion.glu_limit;
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
use_scale = fusion.x_scale != nullptr;
use_gate_scale = fusion.gate_scale != nullptr && use_gate;
x_scale = (const float *) fusion.x_scale;
@@ -651,7 +658,7 @@ static __global__ void mul_mat_vec_q(
gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x];
}
}
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
if (use_scale) {
x_scales = x_scale[ids ? channel_x : 0];
}
@@ -680,11 +687,11 @@ static __global__ void mul_mat_vec_q(
#pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
tmp[j][i] += vec_dot_q_cuda(
vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx_stride*kbx, kqs);
if constexpr (has_fusion) {
if (use_gate) {
tmp_gate[j][i] += vec_dot_q_cuda(
vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx_stride*kbx, kqs);
}
}
}
@@ -739,13 +746,13 @@ static __global__ void mul_mat_vec_q(
if (threadIdx.x == i && (rows_per_cuda_block == 1 || uint32_t(row0 + i) < stride_col_dst)) {
float result = tmp[j][i];
if constexpr (has_fusion) {
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
result *= x_scales;
}
result += x_biases[j];
if (use_gate) {
float gate_value = tmp_gate[j][i];
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
gate_value *= gate_scales;
}
gate_value += gate_biases[j];
@@ -776,7 +783,7 @@ static __global__ void mul_mat_vec_q(
if constexpr (!has_fusion) {
GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, use_scale, use_gate_scale, active_glu, glu_limit, gate_bias, x_bias, x_scale, gate_scale, tmp_gate);
}
if constexpr (type != GGML_TYPE_NVFP4) {
if constexpr (!is_scaled_low_precision_type(type)) {
GGML_UNUSED_VARS(use_scale, use_gate_scale, x_scale, gate_scale, x_scales, gate_scales);
}
}
@@ -802,6 +809,7 @@ static __global__ void mul_mat_vec_q_moe(
constexpr int qk = ggml_cuda_type_traits<type>::qk;
constexpr int qi = ggml_cuda_type_traits<type>::qi;
constexpr int vdr = get_vdr_mmvq(type);
constexpr int kbx_stride = type == GGML_TYPE_F8_E4M3 ? QK8_1 : 1;
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type);
@@ -823,7 +831,7 @@ static __global__ void mul_mat_vec_q_moe(
gate_bias = (const float *) fusion.gate_bias;
active_glu = fusion.glu_op;
glu_limit = fusion.glu_limit;
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
x_scale = (const float *) fusion.x_scale;
gate_scale = (const float *) fusion.gate_scale;
}
@@ -857,10 +865,10 @@ static __global__ void mul_mat_vec_q_moe(
#pragma unroll
for (int i = 0; i < c_rows_per_block; ++i) {
tmp[i] += vec_dot_q_cuda(vx, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs);
tmp[i] += vec_dot_q_cuda(vx, &y[kby], kbx_offset + i*stride_row_x + kbx_stride*kbx, kqs);
if constexpr (has_fusion) {
if (use_gate) {
tmp_gate[i] += vec_dot_q_cuda(vgate, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs);
tmp_gate[i] += vec_dot_q_cuda(vgate, &y[kby], kbx_offset + i*stride_row_x + kbx_stride*kbx, kqs);
}
}
}
@@ -885,7 +893,7 @@ static __global__ void mul_mat_vec_q_moe(
if constexpr (has_fusion) {
const uint32_t bias_idx = channel_x*stride_channel_dst + row0 + threadIdx.x;
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
if (x_scale) {
result *= x_scale[channel_x];
}
@@ -895,7 +903,7 @@ static __global__ void mul_mat_vec_q_moe(
}
if (use_gate) {
float gate_value = tmp_gate[threadIdx.x];
if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (is_scaled_low_precision_type(type)) {
if (gate_scale) {
gate_value *= gate_scale[channel_x];
}
@@ -927,7 +935,7 @@ static __global__ void mul_mat_vec_q_moe(
if constexpr (!has_fusion) {
GGML_UNUSED_VARS(use_gate, tmp_gate, vgate, x_bias, gate_bias, active_glu, glu_limit, x_scale, gate_scale);
} else if constexpr (type != GGML_TYPE_NVFP4) {
} else if constexpr (!is_scaled_low_precision_type(type)) {
GGML_UNUSED_VARS(x_scale, gate_scale);
}
}
@@ -1019,7 +1027,7 @@ static void mul_mat_vec_q_switch_ncols_dst(
const int nsamples_x, const int nsamples_dst, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst,
const int ids_stride, cudaStream_t stream) {
GGML_ASSERT(ncols_x % ggml_blck_size(type) == 0);
GGML_ASSERT(ncols_x % ggml_cuda_type_traits<type>::qk == 0);
GGML_ASSERT(ncols_dst <= MMVQ_MAX_BATCH_SIZE);
const uint3 nchannels_y_fd = ids ? init_fastdiv_values(nchannels_y) : make_uint3(0, 0, 0);
@@ -1265,6 +1273,12 @@ static void mul_mat_vec_q_switch_type(
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
break;
case GGML_TYPE_F8_E4M3:
mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_F8_E4M3>
(vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
break;
case GGML_TYPE_Q2_K:
mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_Q2_K>
(vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
@@ -1387,9 +1401,9 @@ void ggml_cuda_mul_mat_vec_q(
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
GGML_ASSERT( !ids || dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc));
GGML_ASSERT( ids || dst->ne[1] == 1);
// Scale fusion is only allowed for NVFP4 currently as the cost of checking this at run-time in the prologue is
// Scale fusion is only allowed for scaled low-precision types as the cost of checking this at run-time in the prologue is
// non-negligible for some models such as gpt-oss-20b
GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || src0->type == GGML_TYPE_NVFP4);
GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || is_scaled_low_precision_type(src0->type));
if (fusion->x_bias) {
GGML_ASSERT(fusion->x_bias->type == GGML_TYPE_F32);
ggml/src/ggml-cuda/set-rows.cu
+15 -1
View File
@@ -168,7 +168,11 @@ static __global__ void k_set_rows(const src_t * src0_ptr,
const src_t * src0_row = src0 + i01*s01 + i02*s02 + i03*s03;
dst_t * dst_row_ptr = dst + dst_row*s1 + i02*s2 + i03*s3;
dst_row_ptr[i00] = ggml_cuda_cast<dst_t>(src0_row[i00]);
if constexpr (std::is_same_v<dst_t, ggml_fp8_e4m3_t>) {
dst_row_ptr[i00].bits = ggml_cuda_fp32_to_f8_e4m3(src0_row[i00]);
} else {
dst_row_ptr[i00] = ggml_cuda_cast<dst_t>(src0_row[i00]);
}
GGML_UNUSED(ne10);
GGML_UNUSED(ne11);
@@ -257,6 +261,16 @@ static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * s
nb1, nb2, nb3,
stream
);
} else if (dst->type == GGML_TYPE_F8_E4M3) {
set_rows_cuda(
src0_d, src1_d, (ggml_fp8_e4m3_t *) dst->data,
ne00, ne01, ne02, ne03,
ne10, ne11, ne12, ne13,
nb01, nb02, nb03,
nb10, nb11, nb12,
nb1, nb2, nb3,
stream
);
} else if (dst->type == GGML_TYPE_Q4_0) {
set_rows_cuda_quant<idx_t, block_q4_0, QK4_0, quantize_f32_q4_0_block>(
src0_d, src1_d, (block_q4_0*)dst->data,
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-bf16-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_BF16, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_BF16, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_BF16, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f16-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F16, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F16, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F16, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-bf16.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_BF16);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_BF16);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_BF16);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-f16.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_F16);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_F16);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_F16);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-q4_0.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_0);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_0);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_0);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-q4_1.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_1);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_1);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_Q4_1);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-q5_0.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_0);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_0);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_0);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-q5_1.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_1);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_1);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_Q5_1);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f8_e4m3-q8_0.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F8_E4M3, GGML_TYPE_Q8_0);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_F8_E4M3, GGML_TYPE_Q8_0);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_F8_E4M3, GGML_TYPE_Q8_0);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q4_0-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q4_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q4_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q4_0, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q4_1-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q4_1, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q4_1, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q4_1, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q5_0-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q5_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q5_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q5_0, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q5_1-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q5_1, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q5_1, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q5_1, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q8_0-f8_e4m3.cu
+7
View File
@@ -0,0 +1,7 @@
// This file has been autogenerated by generate_cu_files.py, do not edit manually.
#include "../fattn-vec.cuh"
DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q8_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q8_0, GGML_TYPE_F8_E4M3);
DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q8_0, GGML_TYPE_F8_E4M3);
ggml/src/ggml-cuda/template-instances/generate_cu_files.py
+1 -1
View File
@@ -8,7 +8,7 @@ HEAD_SIZES_KQ = [40, 64, 72, 80, 96, 112, 128, 192, 256, 320, 512, 576]
# DKQ -> DV override for asymmetric head dims.
HEAD_SIZES_V_OVERRIDE = {576: 512, 320: 256, 192: 128}
TYPES_KV = ["GGML_TYPE_F16", "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", "GGML_TYPE_BF16"]
TYPES_KV = ["GGML_TYPE_F16", "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", "GGML_TYPE_BF16", "GGML_TYPE_F8_E4M3"]
SOURCE_FATTN_TILE = """// This file has been autogenerated by generate_cu_files.py, do not edit manually.
ggml/src/ggml-cuda/vecdotq.cuh
+19
View File
@@ -360,6 +360,25 @@ static __device__ __forceinline__ float vec_dot_nvfp4_q8_1(
return sum;
}
#define VDR_F8_E4M3_Q8_1_MMVQ 1
static __device__ __forceinline__ float vec_dot_f8_e4m3_q8_1(
const void * __restrict__ vbq,
const block_q8_1 * __restrict__ bq8_1,
const int32_t & kbx,
const int32_t & iqs) {
const ggml_fp8_e4m3_t * bq8 = (const ggml_fp8_e4m3_t *) vbq + kbx;
const int8_t * q8 = bq8_1->qs + 4*iqs;
float sum = 0.0f;
#pragma unroll
for (int i = 0; i < 4; ++i) {
sum += ggml_cuda_f8_e4m3_to_fp32(bq8[4*iqs + i].bits) * q8[i];
}
return __low2float(bq8_1->ds) * sum;
}
#define VDR_Q2_K_Q8_1_MMVQ 1
#define VDR_Q2_K_Q8_1_MMQ 4
ggml/src/ggml-cuda/vendors/cuda.h
Vendored
+1
View File
@@ -3,6 +3,7 @@
#include <cuda_runtime.h>
#include <cuda.h>
#include <cublas_v2.h>
#include <cublasLt.h>
#include <cuda_bf16.h>
#include <cuda_fp16.h>