From a0818450aca6ae0bd3012dbe80d4d602b594abee Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 04:37:57 -0700 Subject: [PATCH 01/23] Blackwell NVFP4 MMQ Kernel --- ggml/src/ggml-cuda/common.cuh | 16 ++++ ggml/src/ggml-cuda/mma.cuh | 26 ++++++ ggml/src/ggml-cuda/mmq.cu | 12 ++- ggml/src/ggml-cuda/mmq.cuh | 152 +++++++++++++++++++++++++++++++- ggml/src/ggml-cuda/mmvq.cu | 3 + ggml/src/ggml-cuda/quantize.cu | 107 ++++++++++++++++++++++ ggml/src/ggml-cuda/quantize.cuh | 5 ++ tests/test-backend-ops.cpp | 4 +- 8 files changed, 320 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 3aec1742ee1a..0c5423121926 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -830,6 +830,22 @@ 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__ 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)) { + return 0; + } + if (x > 448.0f) { + x = 448.0f; + } + + const __nv_fp8_e4m3 xf(x); + return xf.__x; +#else + NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell +#endif // defined(BLACKWELL_MMA_AVAILABLE) +} + __device__ __forceinline__ uint8_t ggml_cuda_float_to_fp4_e2m1(float x, float e) { const uint8_t sign_bit = (x < 0.0f) << 3; float ax = fabsf(x) * e; diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh index b0f674635f1a..a54a3b7b308e 100644 --- a/ggml/src/ggml-cuda/mma.cuh +++ b/ggml/src/ggml-cuda/mma.cuh @@ -1036,6 +1036,32 @@ namespace ggml_cuda_mma { #endif // BLACKWELL_MMA_AVAILABLE } + static __device__ __forceinline__ void mma_nvfp4_blackwell( + tile<16, 8, float> & C, + const tile<16, 8, int> & A, + const tile<8, 8, int> & B, + uint32_t a_scale, + uint32_t b_scale) { +#ifdef BLACKWELL_MMA_AVAILABLE + const int * Axi = (const int *) A.x; + const int * Bxi = (const int *) B.x; + float * Cxi = (float *) C.x; + + asm volatile( + "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X" + ".m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 " + "{%0, %1, %2, %3}, " + "{%4, %5, %6, %7}, " + "{%8, %9}, " + "{%0, %1, %2, %3}, " + "%10, {0, 0}, %11, {0, 0};\n" + : "+f"(Cxi[0]), "+f"(Cxi[1]), "+f"(Cxi[2]), "+f"(Cxi[3]) + : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[1]), "r"(a_scale), "r"(b_scale)); +#else + GGML_UNUSED_VARS(C, A, B, a_scale, b_scale); +#endif // BLACKWELL_MMA_AVAILABLE + } + static __device__ __forceinline__ void mma( tile<16, 8, float> & D, const tile<16, 8, half2> & A, const tile<8, 8, half2> & B) { #ifdef TURING_MMA_AVAILABLE diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 3f01ff5bfb08..008200cc4904 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -123,6 +123,7 @@ void ggml_cuda_mul_mat_q( // TODO: tighter pool buffer size vs q8 path const bool use_native_mxfp4 = blackwell_mma_available(cc) && src0->type == GGML_TYPE_MXFP4; + const bool use_native_nvfp4 = blackwell_mma_available(cc) && src0->type == GGML_TYPE_NVFP4; if (!ids) { const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 + @@ -137,7 +138,10 @@ void ggml_cuda_mul_mat_q( static_assert(sizeof(block_fp4_mmq) == 4 * sizeof(block_q8_1)); quantize_mmq_mxfp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); - + } else if (use_native_nvfp4) { + static_assert(sizeof(block_nvfp4_mmq) == 4 * sizeof(block_q8_1)); + quantize_mmq_nvfp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, + ne11, ne12, ne13, stream); } else { quantize_mmq_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); @@ -149,6 +153,8 @@ void ggml_cuda_mul_mat_q( const int64_t s12 = use_native_mxfp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (8 * QK_MXFP4 * sizeof(int)) // block_fp4_mmq holds 256 values (8 blocks of 32) + : use_native_nvfp4 ? + ne11 * ne10_padded * sizeof(block_nvfp4_mmq) / (QK_K * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; @@ -201,6 +207,9 @@ void ggml_cuda_mul_mat_q( if (use_native_mxfp4) { quantize_mmq_mxfp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); + } else if (use_native_nvfp4) { + quantize_mmq_nvfp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, + ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } else { quantize_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); @@ -209,6 +218,7 @@ void ggml_cuda_mul_mat_q( } const int64_t s12 = use_native_mxfp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (8 * QK_MXFP4 * sizeof(int)) : + use_native_nvfp4 ? ne11 * ne10_padded * sizeof(block_nvfp4_mmq) / (QK_K * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index b1a319de9be9..c46983e919f6 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -12,6 +12,7 @@ using namespace ggml_cuda_mma; #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. #define MMQ_ITER_K 256 #define MMQ_ITER_K_MXFP4_FP4 512 +#define MMQ_ITER_K_NVFP4_FP4 512 #define MMQ_NWARPS 8 typedef void (*load_tiles_mmq_t)(const char * __restrict__ x, int * x_tile, const int kbx0, const int i_max, const int stride); @@ -51,9 +52,15 @@ struct block_fp4_mmq { int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte), 8 blocks of 32 values }; +struct block_nvfp4_mmq { // Same byte size as block_q8_1_mmq and block_fp4_mmq + uint32_t sc4_u32[4]; // 16 UE4M3 scales (1 per 1 values), packed into 4 unint32 + uint32_t qs_u32[32]; // 256 NVFP4 values packed as 4-bit pairs (128 payload bytes) +}; + static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_q8_1_mmq) == 4*sizeof(block_q8_1), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_fp4_mmq size"); +static_assert(sizeof(block_nvfp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_nvfp4_mmq size"); static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { switch (type_x) { @@ -143,7 +150,8 @@ static int get_mmq_y_host(const int cc) { static constexpr __device__ int get_iter_k([[maybe_unused]] const ggml_type type) { #if defined(BLACKWELL_MMA_AVAILABLE) - return type == GGML_TYPE_MXFP4 ? MMQ_ITER_K_MXFP4_FP4 : MMQ_ITER_K; + return type == GGML_TYPE_MXFP4 ? MMQ_ITER_K_MXFP4_FP4 : + type == GGML_TYPE_NVFP4 ? MMQ_ITER_K_NVFP4_FP4 : MMQ_ITER_K; #else return MMQ_ITER_K; #endif // defined(BLACKWELL_MMA_AVAILABLE) @@ -934,6 +942,140 @@ static __device__ __forceinline__ void load_tiles_mxfp4_fp4(const char * __restr } } +#ifdef BLACKWELL_MMA_AVAILABLE +template +static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __restrict__ x, + int * __restrict__ x_tile, + const int kbx0, + const int i_max, + const int stride) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int iter_k = get_iter_k(GGML_TYPE_NVFP4); + constexpr int threads_per_row = iter_k / QK_NVFP4; // each thread processes 1 block + constexpr int rows_per_warp = warp_size / threads_per_row; + + uint32_t * x_u32 = (uint32_t *) x_tile; + + const int txi = threadIdx.x; + const int kbx = txi % threads_per_row; + const int row_in_warp = txi / threads_per_row; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += rows_per_warp * nwarps) { + int i = i0 + threadIdx.y * rows_per_warp + row_in_warp; + + if constexpr (need_check) { + i = min(i, i_max); + } + + const block_nvfp4 * bxi = (const block_nvfp4 *) x + kbx0 + i * stride + kbx; + const int row_base = i * MMQ_MMA_TILE_X_K_NVFP4; + const int q_base = row_base + 8 * kbx; + +#pragma unroll + for (int sub = 0; sub < QK_NVFP4 / QK_NVFP4_SUB; ++sub) { + const uint8_t * p = bxi->qs + 8 * sub; + + x_u32[q_base + 2 * sub + 0] = + (((uint32_t) (p[0] & 0x0F)) << 0) | + (((uint32_t) (p[1] & 0x0F)) << 4) | + (((uint32_t) (p[2] & 0x0F)) << 8) | + (((uint32_t) (p[3] & 0x0F)) << 12) | + (((uint32_t) (p[4] & 0x0F)) << 16) | + (((uint32_t) (p[5] & 0x0F)) << 20) | + (((uint32_t) (p[6] & 0x0F)) << 24) | + (((uint32_t) (p[7] & 0x0F)) << 28); + + x_u32[q_base + 2 * sub + 1] = + (((uint32_t) (p[0] >> 4)) << 0) | + (((uint32_t) (p[1] >> 4)) << 4) | + (((uint32_t) (p[2] >> 4)) << 8) | + (((uint32_t) (p[3] >> 4)) << 12) | + (((uint32_t) (p[4] >> 4)) << 16) | + (((uint32_t) (p[5] >> 4)) << 20) | + (((uint32_t) (p[6] >> 4)) << 24) | + (((uint32_t) (p[7] >> 4)) << 28); + } + + x_u32[i * MMQ_MMA_TILE_X_K_NVFP4 + 64 + kbx] = get_int_b1(bxi->d, 0); + } +} + +template +static __device__ __forceinline__ void vec_dot_nvfp4_nvfp4_mma(const int * __restrict__ x, + const int * __restrict__ y, + float * __restrict__ sum, + const int k00) { + typedef tile<16, 8, int> tile_A; + typedef tile<8, 8, int> tile_B; + typedef tile<16, 8, float> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = 2 * granularity; + constexpr int ntx = rows_per_warp / tile_C::I; + constexpr int nfrags = MMQ_TILE_NE_K / tile_A::J; + constexpr int rows_per_slab = mmq_get_nwarps_device() * tile_C::I; + constexpr int groups_per_slab = mmq_x / tile_C::J; + + const int * x_qs = (const int *) x; + const uint32_t * x_sc = (const uint32_t *) (x_qs + 2 * MMQ_TILE_NE_K); + const int * y_qs = (const int *) y + 4; + const uint32_t * y_sc = (const uint32_t *) y; + const int ty = threadIdx.y; + const int ty_ntx_mod = ty % ntx; + const int ty_ntx_div = ty / ntx; + const int tidx = threadIdx.x / 4 + (threadIdx.x % 2) * 8; + +#pragma unroll + for (int slab_row0 = 0; slab_row0 < mmq_y; slab_row0 += rows_per_slab) { + tile_A A[ntx][nfrags]; + uint32_t scaleA[ntx][nfrags]; + const int i0 = slab_row0 + ty_ntx_div * rows_per_warp; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { +#pragma unroll + for (int frag = 0; frag < nfrags; ++frag) { + const int k0 = k00 + frag * tile_A::J; + load_ldmatrix(A[n][frag], x_qs + (i0 + n * tile_A::I) * MMQ_MMA_TILE_X_K_NVFP4 + k0, MMQ_MMA_TILE_X_K_NVFP4); + scaleA[n][frag] = x_sc[(i0 + n * tile_A::I + tidx) * MMQ_MMA_TILE_X_K_NVFP4 + k0 / tile_A::J]; + } + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx * tile_C::J) { + tile_B B[nfrags]; + uint32_t scaleB[nfrags]; + const int sum_j = slab_row0 / rows_per_slab * groups_per_slab + j0 / tile_C::J; + +#pragma unroll + for (int frag = 0; frag < nfrags; ++frag) { + const int k0 = frag * tile_B::J; + load_generic(B[frag], y_qs + (ty_ntx_mod * tile_C::J + j0) * MMQ_TILE_Y_K + k0, MMQ_TILE_Y_K); + scaleB[frag] = y_sc[(ty_ntx_mod * tile_C::J + j0 + tidx) * MMQ_TILE_Y_K + frag]; + } + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + float * sum_n = sum + (sum_j + n) * tile_C::ne; + +#pragma unroll + for (int frag = 0; frag < nfrags; ++frag) { + tile_C C = {}; + mma_nvfp4_blackwell(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + sum_n[l] += C.x[l]; + } + } + } + } + } +} +#endif // BLACKWELL_MMA_AVAILABLE + template static __device__ __forceinline__ void load_tiles_nvfp4(const char * __restrict__ x, @@ -3270,8 +3412,13 @@ struct mmq_type_traits { template struct mmq_type_traits { static constexpr int vdr = VDR_NVFP4_Q8_1_MMQ; +#ifdef BLACKWELL_MMA_AVAILABLE + static constexpr load_tiles_mmq_t load_tiles = load_tiles_nvfp4_nvfp4; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_nvfp4_nvfp4_mma; +#else static constexpr load_tiles_mmq_t load_tiles = load_tiles_nvfp4; static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_16_q8_1_mma; +#endif // BLACKWELL_MMA_AVAILABLE static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_16_q8_1_dp4a; }; @@ -3406,7 +3553,8 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( #if defined(BLACKWELL_MMA_AVAILABLE) // FP4 tile stores 8 blocks - constexpr int ne_block = (type == GGML_TYPE_MXFP4) ? 8 * QK_MXFP4 : 4 * QK8_1; + constexpr int ne_block = (type == GGML_TYPE_MXFP4) ? 8 * QK_MXFP4 : + (type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1; #else constexpr int ne_block = 4 * QK8_1; #endif // defined(BLACKWELL_MMA_AVAILABLE) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 8f55cace1a1e..da48f313a38b 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -115,6 +115,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_pascal_older(gg case GGML_TYPE_IQ4_NL: return 6; case GGML_TYPE_IQ4_XS: return 5; case GGML_TYPE_MXFP4: return 4; + case GGML_TYPE_NVFP4: return 4; case GGML_TYPE_Q2_K: return 4; case GGML_TYPE_Q3_K: return 4; case GGML_TYPE_Q4_0: return 6; @@ -135,6 +136,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_turing_plus(ggm case GGML_TYPE_IQ3_S: return 6; case GGML_TYPE_IQ3_XXS: return 7; case GGML_TYPE_MXFP4: return 7; + case GGML_TYPE_NVFP4: return 8; case GGML_TYPE_Q2_K: return 7; case GGML_TYPE_Q3_K: return 5; default: return MMVQ_MAX_BATCH_SIZE; @@ -221,6 +223,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type case GGML_TYPE_IQ4_NL: return 7; case GGML_TYPE_IQ4_XS: return 5; case GGML_TYPE_MXFP4: return 5; + case GGML_TYPE_NVFP4: return 5; case GGML_TYPE_Q3_K: return 4; case GGML_TYPE_Q4_0: return 7; case GGML_TYPE_Q4_1: return 7; diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 4300ffc148cf..7b8180055c76 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -70,6 +70,92 @@ __device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) { return static_cast(biased); } +template +static __global__ void quantize_mmq_nvfp4( + const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, + const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2) { + const int lane_id = threadIdx.x & 31; + + const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * 8; + if (i0_base >= ne0) { + return; + } + + const int64_t i1 = blockIdx.x; + const int64_t i2 = blockIdx.z % ne2; + const int64_t i3 = blockIdx.z / ne2; + const int64_t i01 = has_ids ? ids[i1] : i1; + const int64_t k_block = i0_base / QK_K; + const int64_t blocks_per_col = (ne0 + QK_K - 1) / QK_K; + if (k_block >= blocks_per_col) { + return; + } + + const int64_t ib = blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x; + block_nvfp4_mmq * y = (block_nvfp4_mmq *) vy; + block_nvfp4_mmq * yb = y + ib; + + float vals_raw[8]; + float amax_raw = 0.0f; + const int64_t base_idx = i3 * s03 + i2 * s02 + i01 * s01; +#pragma unroll + for (int k = 0; k < 8; k++) { + const int64_t i00 = i0_base + k; + if (i00 < ne00) { + const float v = x[base_idx + i00]; + vals_raw[k] = v; + amax_raw = fmaxf(amax_raw, fabsf(v)); + } else { + vals_raw[k] = 0.0f; + } + } + + const float sub_max = fmaxf(amax_raw, __shfl_xor_sync(0xFFFFFFFFu, amax_raw, 1)); + static constexpr int test_offsets[5] = { 0, -1, 1, -2, 2}; + const int first_fp8_code = (int) ggml_cuda_fp32_to_ue4m3(sub_max * (1.0f / 6.0f)); + + float best_err = FLT_MAX; + uint8_t fp8_code = 0; + float subblock_scale = 0.0f; + +#pragma unroll // Check +/- 2 to find best code to reduce NVFP4 activation loss. Negligible overhead on Blackwell. + for (int i = 0; i < 5; i++) { + const int test_code = first_fp8_code + test_offsets[i]; + const uint8_t code = (uint8_t) test_code; + const float test_scale = ggml_cuda_ue4m3_to_fp32(code); + const float test_inv_scale = test_scale > 0.0f ? 0.5f / test_scale : 0.0f; + float cur_err = 0.0f; +#pragma unroll + for (int k = 0; k < 8; ++k) { + const float v = vals_raw[k]; + const uint8_t q = ggml_cuda_float_to_fp4_e2m1(v, test_inv_scale); + const float err_diff = fabsf(v) - fabsf(kvalues_mxfp4[q & 0x7]) * test_scale; + cur_err = fmaf(err_diff, err_diff, cur_err); + } + + cur_err = cur_err + __shfl_xor_sync(__activemask(), cur_err, 1); + + if (cur_err < best_err) { + best_err = cur_err; + fp8_code = test_code; + subblock_scale = test_scale; + } + } + + const float inv_scale = subblock_scale > 0.0f ? 0.5f / subblock_scale : 0.0f; + uint32_t q_word = 0; +#pragma unroll // this is faster than the previous __nv_fp4x4_e2m1 + for (int k = 0; k < 8; ++k) { + q_word |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k], inv_scale) << (4 * k); + } + yb->qs_u32[lane_id] = q_word; + + if ((lane_id & 1) == 0) { + reinterpret_cast(yb->sc4_u32)[lane_id >> 1] = fp8_code; + } +} + // quantize values in the format mxfp4 is stored which is interleaved nibbles // i.e. a block a0-a31 is represented as a0a16,a1a17 ...a15a31 static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, @@ -341,3 +427,24 @@ void quantize_mmq_mxfp4_cuda(const float * x, quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); } + +void quantize_mmq_nvfp4_cuda( + const float * x, const int32_t * ids, void * vy, const ggml_type type_src0, + const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { + GGML_ASSERT(type_src0 == GGML_TYPE_NVFP4); + GGML_ASSERT(ne00 % 8 == 0); + GGML_ASSERT(ne0 > 0); + + constexpr int nvfp4_block_size = 128; + const int64_t block_num_y = (ne0 + 8 * nvfp4_block_size - 1) / (8 * nvfp4_block_size); + const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); + const dim3 block_size(nvfp4_block_size, 1, 1); + if (ids) { + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + } else { + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + } +} diff --git a/ggml/src/ggml-cuda/quantize.cuh b/ggml/src/ggml-cuda/quantize.cuh index 6a91df635788..8558b45ceb3a 100644 --- a/ggml/src/ggml-cuda/quantize.cuh +++ b/ggml/src/ggml-cuda/quantize.cuh @@ -39,3 +39,8 @@ void quantize_mmq_mxfp4_cuda(const float * x, int64_t ne2, int64_t ne3, cudaStream_t stream); + +void quantize_mmq_nvfp4_cuda( + const float * x, const int32_t * ids, void * vy, + ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03, + int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 828a9c14a45d..474c1269768b 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -3765,7 +3765,7 @@ struct test_mul_mat : public test_case { double max_nmse_err(ggml_backend_t backend) override { // for blackwell we quantize activations to mxfp4 instead of q8_1 so we add higher tolerance - if (type_a == GGML_TYPE_MXFP4 && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { + if ((type_a == GGML_TYPE_MXFP4 || type_a == GGML_TYPE_NVFP4) && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { return 2e-2; } return max_nmse_err(); @@ -3901,7 +3901,7 @@ struct test_mul_mat_id : public test_case { double max_nmse_err(ggml_backend_t backend) override { // for blackwell we quantize activations to mxfp4 instead of q8_1 so we add higher tolerance - if (type_a == GGML_TYPE_MXFP4 && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { + if ((type_a == GGML_TYPE_MXFP4 || type_a == GGML_TYPE_NVFP4) && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { return 2e-2; } return max_nmse_err(); From 9fb7e84034359545afb322e84d8fbeb43967e0ef Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 15:53:31 -0700 Subject: [PATCH 02/23] Removed whitespace --- ggml/src/ggml-cuda/common.cuh | 4 ++-- ggml/src/ggml-cuda/quantize.cu | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 0c5423121926..2ad262fa1c28 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -836,13 +836,13 @@ static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) { return 0; } if (x > 448.0f) { - x = 448.0f; + x = F; } const __nv_fp8_e4m3 xf(x); return xf.__x; #else - NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell + NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell #endif // defined(BLACKWELL_MMA_AVAILABLE) } diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 7b8180055c76..19cd29fff5b3 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -120,7 +120,7 @@ static __global__ void quantize_mmq_nvfp4( float subblock_scale = 0.0f; #pragma unroll // Check +/- 2 to find best code to reduce NVFP4 activation loss. Negligible overhead on Blackwell. - for (int i = 0; i < 5; i++) { + for (int i = 0; i < 5; i++) { const int test_code = first_fp8_code + test_offsets[i]; const uint8_t code = (uint8_t) test_code; const float test_scale = ggml_cuda_ue4m3_to_fp32(code); @@ -136,7 +136,7 @@ static __global__ void quantize_mmq_nvfp4( cur_err = cur_err + __shfl_xor_sync(__activemask(), cur_err, 1); - if (cur_err < best_err) { + if (cur_err < best_err) { best_err = cur_err; fp8_code = test_code; subblock_scale = test_scale; From 0bcf7b29b86b67bf2142acb5faca27fcfe204922 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 15:57:31 -0700 Subject: [PATCH 03/23] Added FP8 Max definition and description --- ggml/src/ggml-cuda/common.cuh | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 2ad262fa1c28..61220a6ec8f4 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -830,13 +830,14 @@ 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 } +#define UE4M3_MAX 448.0f // The maximum representable value for Ue4m3 is 448.0f for NVFP4 FP8 scales 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)) { return 0; } - if (x > 448.0f) { - x = F; + if (x > UE4M3_MAX) { + x = UE4M3_MAX; } const __nv_fp8_e4m3 xf(x); From 4625a7cc4746c45467be5c1306f058c5291ed563 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 16:02:29 -0700 Subject: [PATCH 04/23] Fixed 'f' typo --- ggml/src/ggml-cuda/common.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 61220a6ec8f4..6231ea4fb1b3 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -843,7 +843,7 @@ static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) { const __nv_fp8_e4m3 xf(x); return xf.__x; #else - NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell + NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell #endif // defined(BLACKWELL_MMA_AVAILABLE) } From 3ea6b59d7b9b79df0f390e5767ee92d2bf8c8dfd Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 16:03:41 -0700 Subject: [PATCH 05/23] Removed whitespace from comment --- ggml/src/ggml-cuda/common.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 6231ea4fb1b3..61220a6ec8f4 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -843,7 +843,7 @@ static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) { const __nv_fp8_e4m3 xf(x); return xf.__x; #else - NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell + NO_DEVICE_CODE; // Used only for NVFP4 Scales for Activations, only for Blackwell #endif // defined(BLACKWELL_MMA_AVAILABLE) } From db5957e7cfc7b14a9a66ad83683c12f019474786 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 21:55:40 -0700 Subject: [PATCH 06/23] Guard Blackwell NVFP4 quantizer for Blackwell only --- ggml/src/ggml-cuda/quantize.cu | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 19cd29fff5b3..65d9298a407a 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -75,6 +75,8 @@ static __global__ void quantize_mmq_nvfp4( const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2) { +#if defined(BLACKWELL_MMA_AVAILABLE) + const int lane_id = threadIdx.x & 31; const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * 8; @@ -154,6 +156,10 @@ static __global__ void quantize_mmq_nvfp4( if ((lane_id & 1) == 0) { reinterpret_cast(yb->sc4_u32)[lane_id >> 1] = fp8_code; } +#else + NO_DEVICE_CODE; // This is for Blackwell NVFP4 activations only. +#endif // defined(BLACKWELL_MMA_AVAILABLE) + } // quantize values in the format mxfp4 is stored which is interleaved nibbles From 83b412f0727136d60a65e8001ebae6d111624b59 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Tue, 14 Apr 2026 23:19:26 -0700 Subject: [PATCH 07/23] Merged vec_dot_fp4_fp4_mma together --- ggml/src/ggml-cuda/mmq.cuh | 164 +++++++++++-------------------------- 1 file changed, 49 insertions(+), 115 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index c46983e919f6..ef3986c0e2ce 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -1002,73 +1002,78 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res } } -template -static __device__ __forceinline__ void vec_dot_nvfp4_nvfp4_mma(const int * __restrict__ x, - const int * __restrict__ y, - float * __restrict__ sum, - const int k00) { +// Shared MMA kernel for MXFP4 and NVFP4 on Blackwell. +// Both quantizations encode values as e2m1 (FP4) and produce one uint32 scale per +// m16n8k64 MMA call; only the PTX kind (scale_vec::2X ue8m0 vs scale_vec::4X ue4m3) +// and the per-type stride constant differ. +template +static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restrict__ x, + const int * __restrict__ y, + float * __restrict__ sum, + const int k00) { + static_assert(type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4, + "vec_dot_fp4_fp4_mma: type must be MXFP4 or NVFP4"); + typedef tile<16, 8, int> tile_A; typedef tile<8, 8, int> tile_B; typedef tile<16, 8, float> tile_C; - constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int stride = mmq_get_mma_tile_x_k(type); + constexpr int granularity = mmq_get_granularity_device(mmq_x); constexpr int rows_per_warp = 2 * granularity; - constexpr int ntx = rows_per_warp / tile_C::I; - constexpr int nfrags = MMQ_TILE_NE_K / tile_A::J; - constexpr int rows_per_slab = mmq_get_nwarps_device() * tile_C::I; - constexpr int groups_per_slab = mmq_x / tile_C::J; + constexpr int ntx = rows_per_warp / tile_C::I; + constexpr int nfrags = MMQ_TILE_NE_K / tile_A::J; + + y += (threadIdx.y % ntx) * (tile_C::J * MMQ_TILE_Y_K); const int * x_qs = (const int *) x; const uint32_t * x_sc = (const uint32_t *) (x_qs + 2 * MMQ_TILE_NE_K); const int * y_qs = (const int *) y + 4; const uint32_t * y_sc = (const uint32_t *) y; - const int ty = threadIdx.y; - const int ty_ntx_mod = ty % ntx; - const int ty_ntx_div = ty / ntx; + + // 2 threads per quad supply the packed scale register to the block_scale MMA, + // see https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-block-scaling const int tidx = threadIdx.x / 4 + (threadIdx.x % 2) * 8; + const int i0 = (threadIdx.y / ntx) * rows_per_warp; -#pragma unroll - for (int slab_row0 = 0; slab_row0 < mmq_y; slab_row0 += rows_per_slab) { - tile_A A[ntx][nfrags]; - uint32_t scaleA[ntx][nfrags]; - const int i0 = slab_row0 + ty_ntx_div * rows_per_warp; + tile_A A[ntx][nfrags]; + uint32_t scaleA[ntx][nfrags]; #pragma unroll - for (int n = 0; n < ntx; ++n) { + for (int n = 0; n < ntx; ++n) { #pragma unroll - for (int frag = 0; frag < nfrags; ++frag) { - const int k0 = k00 + frag * tile_A::J; - load_ldmatrix(A[n][frag], x_qs + (i0 + n * tile_A::I) * MMQ_MMA_TILE_X_K_NVFP4 + k0, MMQ_MMA_TILE_X_K_NVFP4); - scaleA[n][frag] = x_sc[(i0 + n * tile_A::I + tidx) * MMQ_MMA_TILE_X_K_NVFP4 + k0 / tile_A::J]; - } + for (int frag = 0; frag < nfrags; ++frag) { + const int k0 = k00 + frag * tile_A::J; + load_ldmatrix(A[n][frag], x_qs + (i0 + n * tile_A::I) * stride + k0, stride); + scaleA[n][frag] = x_sc[(i0 + n * tile_A::I + tidx) * stride + k0 / tile_A::J]; } + } #pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += ntx * tile_C::J) { - tile_B B[nfrags]; - uint32_t scaleB[nfrags]; - const int sum_j = slab_row0 / rows_per_slab * groups_per_slab + j0 / tile_C::J; + for (int j0 = 0; j0 < mmq_x; j0 += ntx * tile_C::J) { + tile_B B[nfrags]; + uint32_t scaleB[nfrags]; #pragma unroll - for (int frag = 0; frag < nfrags; ++frag) { - const int k0 = frag * tile_B::J; - load_generic(B[frag], y_qs + (ty_ntx_mod * tile_C::J + j0) * MMQ_TILE_Y_K + k0, MMQ_TILE_Y_K); - scaleB[frag] = y_sc[(ty_ntx_mod * tile_C::J + j0 + tidx) * MMQ_TILE_Y_K + frag]; - } + for (int frag = 0; frag < nfrags; ++frag) { + const int k0 = frag * tile_B::J; + load_generic(B[frag], y_qs + j0 * MMQ_TILE_Y_K + k0, MMQ_TILE_Y_K); + scaleB[frag] = y_sc[(j0 + tidx) * MMQ_TILE_Y_K + frag]; + } #pragma unroll - for (int n = 0; n < ntx; ++n) { - float * sum_n = sum + (sum_j + n) * tile_C::ne; - + for (int n = 0; n < ntx; ++n) { #pragma unroll - for (int frag = 0; frag < nfrags; ++frag) { - tile_C C = {}; + for (int frag = 0; frag < nfrags; ++frag) { + tile_C C = {}; + if constexpr (type == GGML_TYPE_MXFP4) { + mma_block_scaled (C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); + } else { mma_nvfp4_blackwell(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); - + } #pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum_n[l] += C.x[l]; - } + for (int l = 0; l < tile_C::ne; ++l) { + sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; } } } @@ -1305,77 +1310,6 @@ static __device__ __forceinline__ void vec_dot_q8_0_q8_1_mma( #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } -template -static __device__ __forceinline__ void vec_dot_mxfp4_mxfp4_mma(const int * __restrict__ x, - const int * __restrict__ y, - float * __restrict__ sum, - const int k00) { - typedef tile<16, 8, int> tile_A; - typedef tile<8, 8, int> tile_B; - typedef tile<16, 8, float> tile_C; // Output is float for native scaled MMA - - constexpr int granularity = mmq_get_granularity_device(mmq_x); - constexpr int rows_per_warp = 2 * granularity; - constexpr int ntx = rows_per_warp / tile_C::I; // Number of x minitiles per warp. - - y += (threadIdx.y % ntx) * (tile_C::J * MMQ_TILE_Y_FP4_K); - - // Match layout from load_tiles_mxfp4_fp4 - const int * x_qs = (const int *) x; - const uint32_t * x_sc = (const uint32_t *) (x_qs + 2 * MMQ_TILE_NE_K); - const int * y_qs = (const int *) y + 4; - const uint32_t * y_sc = (const uint32_t *) y; - - // tile_A has a length of 64 logical values vs. 32 values in block_mxfp4 - tile_A A[ntx][MMQ_TILE_NE_K / (2 * QI_MXFP4)]; - uint32_t scaleA[ntx][MMQ_TILE_NE_K / (2 * QI_MXFP4)]; - - // Block scale - // Each thread has to point to a 4 byte scale value - // https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-block-scaling - - const int i0 = (threadIdx.y / ntx) * rows_per_warp; - -#pragma unroll - for (int n = 0; n < ntx; ++n) { -#pragma unroll - for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 2 * QI_MXFP4) { - const int k0 = k00 + k01; - - load_ldmatrix(A[n][k01 / (2 * QI_MXFP4)], x_qs + (i0 + n * tile_A::I) * MMQ_MMA_TILE_X_K_FP4 + k0, - MMQ_MMA_TILE_X_K_FP4); - - // based on block-scaling document, 2 threads in each quad need to supply to the scale value - const int tidx = threadIdx.x / 4 + (threadIdx.x % 2) * 8; - scaleA[n][k01 / (2 * QI_MXFP4)] = - *(x_sc + (i0 + n * tile_A::I + tidx) * MMQ_MMA_TILE_X_K_FP4 + k0 / (2 * QI_MXFP4)); - } - } - -#pragma unroll - for (int j0 = 0; j0 < mmq_x; j0 += ntx * tile_C::J) { -#pragma unroll - for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 2 * QI_MXFP4) { - tile_B B; - uint32_t scaleB; // 2xN scales - - load_generic(B, y_qs + j0 * MMQ_TILE_Y_FP4_K + k01, MMQ_TILE_Y_FP4_K); - - scaleB = y_sc[(j0 + threadIdx.x / 4) * MMQ_TILE_Y_FP4_K + k01 / (2 * QI_MXFP4)]; - -#pragma unroll - for (int n = 0; n < ntx; ++n) { - tile_C C; - - mma_block_scaled(C, A[n][k01 / (2 * QI_MXFP4)], B, scaleA[n][k01 / (2 * QI_MXFP4)], scaleB); -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; - } - } - } - } -} template static __device__ __forceinline__ void vec_dot_q8_1_q8_1_dp4a( @@ -3401,7 +3335,7 @@ struct mmq_type_traits { static constexpr int vdr = VDR_MXFP4_Q8_1_MMQ; #ifdef BLACKWELL_MMA_AVAILABLE static constexpr load_tiles_mmq_t load_tiles = load_tiles_mxfp4_fp4; - static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_mxfp4_mxfp4_mma; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_fp4_fp4_mma; #else static constexpr load_tiles_mmq_t load_tiles = load_tiles_mxfp4; static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; @@ -3414,7 +3348,7 @@ struct mmq_type_traits { static constexpr int vdr = VDR_NVFP4_Q8_1_MMQ; #ifdef BLACKWELL_MMA_AVAILABLE static constexpr load_tiles_mmq_t load_tiles = load_tiles_nvfp4_nvfp4; - static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_nvfp4_nvfp4_mma; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_fp4_fp4_mma; #else static constexpr load_tiles_mmq_t load_tiles = load_tiles_nvfp4; static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_16_q8_1_mma; From c3188065c7162106c9c65a3f217bdafcaaf053b6 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 09:10:35 -0700 Subject: [PATCH 08/23] Refactored to use 76-byte MMQ_MMA_TILE_X_K_FP4 and block_fp4_mmq instead of block_nvfp4, removed UE4M3 max cap check, merged use_native_mxfp4/nvfp4 into use_native_fp4, merged quantize_mmq_nvfp4/mxfp4/cuda to quantize_mmq_fp4_Cuda, merged mma/mxfp4/nvfp4 into one templated mma_block_scaled_fp4 --- ggml/src/ggml-cuda/common.cuh | 5 -- ggml/src/ggml-cuda/mma.cuh | 58 +++++++----------- ggml/src/ggml-cuda/mmq.cu | 32 ++++------ ggml/src/ggml-cuda/mmq.cuh | 75 ++++++++--------------- ggml/src/ggml-cuda/quantize.cu | 103 ++++++++++++++------------------ ggml/src/ggml-cuda/quantize.cuh | 7 +-- 6 files changed, 103 insertions(+), 177 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 61220a6ec8f4..10817505d9f0 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -830,16 +830,11 @@ 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 } -#define UE4M3_MAX 448.0f // The maximum representable value for Ue4m3 is 448.0f for NVFP4 FP8 scales 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)) { return 0; } - if (x > UE4M3_MAX) { - x = UE4M3_MAX; - } - const __nv_fp8_e4m3 xf(x); return xf.__x; #else diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh index a54a3b7b308e..79bb2934c5f9 100644 --- a/ggml/src/ggml-cuda/mma.cuh +++ b/ggml/src/ggml-cuda/mma.cuh @@ -1015,50 +1015,34 @@ namespace ggml_cuda_mma { #endif // AMD_MFMA_AVAILABLE } - static __device__ __forceinline__ void mma_block_scaled(tile<16, 8, float> & D, - const tile<16, 8, int> & A, - const tile<8, 8, int> & B, - uint32_t a_scale, - uint32_t b_scale) { + template + static __device__ __forceinline__ void mma_block_scaled_fp4(tile<16, 8, float> & D, + const tile<16, 8, int> & A, + const tile<8, 8, int> & B, + uint32_t a_scale, + uint32_t b_scale) { #ifdef BLACKWELL_MMA_AVAILABLE const int * Axi = (const int *) A.x; const int * Bxi = (const int *) B.x; float * Dxi = (float *) D.x; - asm volatile( - "mma.sync.aligned.kind::mxf4.block_scale.scale_vec::2X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue8m0 " - "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3}, " - "%10, {0, 0}, %11, {0, 0};" - : "+f"(Dxi[0]), "+f"(Dxi[1]), "+f"(Dxi[2]), "+f"(Dxi[3]) - : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[1]), "r"(a_scale), "r"(b_scale)); + if constexpr (type == GGML_TYPE_MXFP4) { + asm volatile( + "mma.sync.aligned.kind::mxf4.block_scale.scale_vec::2X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue8m0 " + "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3}, " + "%10, {0, 0}, %11, {0, 0};" + : "+f"(Dxi[0]), "+f"(Dxi[1]), "+f"(Dxi[2]), "+f"(Dxi[3]) + : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[1]), "r"(a_scale), "r"(b_scale)); + } else { + asm volatile( + "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 " + "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3}, " + "%10, {0, 0}, %11, {0, 0};" + : "+f"(Dxi[0]), "+f"(Dxi[1]), "+f"(Dxi[2]), "+f"(Dxi[3]) + : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[1]), "r"(a_scale), "r"(b_scale)); + } #else GGML_UNUSED_VARS(D, A, B, a_scale, b_scale); -#endif // BLACKWELL_MMA_AVAILABLE - } - - static __device__ __forceinline__ void mma_nvfp4_blackwell( - tile<16, 8, float> & C, - const tile<16, 8, int> & A, - const tile<8, 8, int> & B, - uint32_t a_scale, - uint32_t b_scale) { -#ifdef BLACKWELL_MMA_AVAILABLE - const int * Axi = (const int *) A.x; - const int * Bxi = (const int *) B.x; - float * Cxi = (float *) C.x; - - asm volatile( - "mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X" - ".m16n8k64.row.col.f32.e2m1.e2m1.f32.ue4m3 " - "{%0, %1, %2, %3}, " - "{%4, %5, %6, %7}, " - "{%8, %9}, " - "{%0, %1, %2, %3}, " - "%10, {0, 0}, %11, {0, 0};\n" - : "+f"(Cxi[0]), "+f"(Cxi[1]), "+f"(Cxi[2]), "+f"(Cxi[3]) - : "r"(Axi[0]), "r"(Axi[1]), "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[0]), "r"(Bxi[1]), "r"(a_scale), "r"(b_scale)); -#else - GGML_UNUSED_VARS(C, A, B, a_scale, b_scale); #endif // BLACKWELL_MMA_AVAILABLE } diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 008200cc4904..0a6f898394aa 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -122,8 +122,7 @@ void ggml_cuda_mul_mat_q( || GGML_CUDA_CC_IS_CDNA(cc); // TODO: tighter pool buffer size vs q8 path - const bool use_native_mxfp4 = blackwell_mma_available(cc) && src0->type == GGML_TYPE_MXFP4; - const bool use_native_nvfp4 = blackwell_mma_available(cc) && src0->type == GGML_TYPE_NVFP4; + const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); if (!ids) { const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 + @@ -134,14 +133,11 @@ void ggml_cuda_mul_mat_q( const int64_t s11 = src1->nb[1] / ts_src1; const int64_t s12 = src1->nb[2] / ts_src1; const int64_t s13 = src1->nb[3] / ts_src1; - if (use_native_mxfp4) { + if (use_native_fp4) { static_assert(sizeof(block_fp4_mmq) == 4 * sizeof(block_q8_1)); - quantize_mmq_mxfp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, - ne11, ne12, ne13, stream); - } else if (use_native_nvfp4) { - static_assert(sizeof(block_nvfp4_mmq) == 4 * sizeof(block_q8_1)); - quantize_mmq_nvfp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, + quantize_mmq_fp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); + } else { quantize_mmq_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); @@ -150,12 +146,8 @@ void ggml_cuda_mul_mat_q( } // Stride depends on quantization format - const int64_t s12 = use_native_mxfp4 ? - ne11 * ne10_padded * sizeof(block_fp4_mmq) / - (8 * QK_MXFP4 * sizeof(int)) // block_fp4_mmq holds 256 values (8 blocks of 32) - : use_native_nvfp4 ? - ne11 * ne10_padded * sizeof(block_nvfp4_mmq) / (QK_K * sizeof(int)) - : + const int64_t s12 = use_native_fp4 ? + ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : // block_fp4_mmq holds 256 values (8 blocks of 32) ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; @@ -204,11 +196,8 @@ void ggml_cuda_mul_mat_q( const int64_t s12 = src1->nb[2] / ts_src1; const int64_t s13 = src1->nb[3] / ts_src1; - if (use_native_mxfp4) { - quantize_mmq_mxfp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, - ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); - } else if (use_native_nvfp4) { - quantize_mmq_nvfp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, + if (use_native_fp4) { + quantize_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } else { quantize_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, @@ -217,9 +206,8 @@ void ggml_cuda_mul_mat_q( CUDA_CHECK(cudaGetLastError()); } - const int64_t s12 = use_native_mxfp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (8 * QK_MXFP4 * sizeof(int)) : - use_native_nvfp4 ? ne11 * ne10_padded * sizeof(block_nvfp4_mmq) / (QK_K * sizeof(int)) : - ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); + const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : + ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; // Note that ne02 is used instead of ne12 because the number of y channels determines the z dimension of the CUDA grid. diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index ef3986c0e2ce..405231d2951a 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -11,8 +11,7 @@ using namespace ggml_cuda_mma; #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. #define MMQ_ITER_K 256 -#define MMQ_ITER_K_MXFP4_FP4 512 -#define MMQ_ITER_K_NVFP4_FP4 512 +#define MMQ_ITER_K_FP4 512 #define MMQ_NWARPS 8 typedef void (*load_tiles_mmq_t)(const char * __restrict__ x, int * x_tile, const int kbx0, const int i_max, const int stride); @@ -47,20 +46,14 @@ struct block_q8_1_mmq { int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each }; -struct block_fp4_mmq { - uint32_t d4[4]; // 8 E8M0 scales (1 per 32 values), 2 packed per uint32: d4[0]={s0,s1}, d4[1]={s2,s3}, etc. +struct block_fp4_mmq { // Used for both MXFP4[8 E8M0 scales (1 per 32)) and NVFP4 (16 UE4M3 scales (1 per 16)) + uint32_t d4[4]; // 2 packed per uint32: d4[0]={s0,s1}, d4[1]={s2,s3}, etc. int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte), 8 blocks of 32 values }; -struct block_nvfp4_mmq { // Same byte size as block_q8_1_mmq and block_fp4_mmq - uint32_t sc4_u32[4]; // 16 UE4M3 scales (1 per 1 values), packed into 4 unint32 - uint32_t qs_u32[32]; // 256 NVFP4 values packed as 4-bit pairs (128 payload bytes) -}; - static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_q8_1_mmq) == 4*sizeof(block_q8_1), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_fp4_mmq size"); -static_assert(sizeof(block_nvfp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_nvfp4_mmq size"); static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { switch (type_x) { @@ -150,8 +143,7 @@ static int get_mmq_y_host(const int cc) { static constexpr __device__ int get_iter_k([[maybe_unused]] const ggml_type type) { #if defined(BLACKWELL_MMA_AVAILABLE) - return type == GGML_TYPE_MXFP4 ? MMQ_ITER_K_MXFP4_FP4 : - type == GGML_TYPE_NVFP4 ? MMQ_ITER_K_NVFP4_FP4 : MMQ_ITER_K; + return MMQ_ITER_K_FP4; #else return MMQ_ITER_K; #endif // defined(BLACKWELL_MMA_AVAILABLE) @@ -221,8 +213,8 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml } #define MMQ_MMA_TILE_X_K_Q8_0 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) -#define MMQ_MMA_TILE_X_K_FP4 (2*MMQ_TILE_NE_K + 8 + 4) // MXFP4 -#define MMQ_MMA_TILE_X_K_NVFP4 (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) // NVFP4 +#define MMQ_MMA_TILE_X_K_FP4 (2*MMQ_TILE_NE_K + 8 + 4) // MXFP4 and NVFP4 Blackwell +#define MMQ_MMA_TILE_X_K_NVFP4 (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) // NVFP4 Generic #define MMQ_MMA_TILE_X_K_Q8_1 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) #define MMQ_MMA_TILE_X_K_Q2_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K + 4) #define MMQ_MMA_TILE_X_K_Q3_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) @@ -248,7 +240,11 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { case GGML_TYPE_Q8_0: return MMQ_MMA_TILE_X_K_Q8_0; // tile sizes are the same for Q8_1 and FP4 for blackwell case GGML_TYPE_MXFP4: return MMQ_MMA_TILE_X_K_Q8_1; - case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; +#if defined(BLACKWELL_MMA_AVAILABLE) + case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_FP4; +#else + case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; +#endif case GGML_TYPE_Q2_K: return MMQ_MMA_TILE_X_K_Q2_K; case GGML_TYPE_Q3_K: return MMQ_MMA_TILE_X_K_Q3_K; case GGML_TYPE_Q4_K: return MMQ_MMA_TILE_X_K_Q8_1; @@ -970,35 +966,18 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res } const block_nvfp4 * bxi = (const block_nvfp4 *) x + kbx0 + i * stride + kbx; - const int row_base = i * MMQ_MMA_TILE_X_K_NVFP4; + const int row_base = i * MMQ_MMA_TILE_X_K_FP4; const int q_base = row_base + 8 * kbx; + const uint32_t * src_qs = reinterpret_cast(bxi->qs); + #pragma unroll for (int sub = 0; sub < QK_NVFP4 / QK_NVFP4_SUB; ++sub) { - const uint8_t * p = bxi->qs + 8 * sub; - - x_u32[q_base + 2 * sub + 0] = - (((uint32_t) (p[0] & 0x0F)) << 0) | - (((uint32_t) (p[1] & 0x0F)) << 4) | - (((uint32_t) (p[2] & 0x0F)) << 8) | - (((uint32_t) (p[3] & 0x0F)) << 12) | - (((uint32_t) (p[4] & 0x0F)) << 16) | - (((uint32_t) (p[5] & 0x0F)) << 20) | - (((uint32_t) (p[6] & 0x0F)) << 24) | - (((uint32_t) (p[7] & 0x0F)) << 28); - - x_u32[q_base + 2 * sub + 1] = - (((uint32_t) (p[0] >> 4)) << 0) | - (((uint32_t) (p[1] >> 4)) << 4) | - (((uint32_t) (p[2] >> 4)) << 8) | - (((uint32_t) (p[3] >> 4)) << 12) | - (((uint32_t) (p[4] >> 4)) << 16) | - (((uint32_t) (p[5] >> 4)) << 20) | - (((uint32_t) (p[6] >> 4)) << 24) | - (((uint32_t) (p[7] >> 4)) << 28); + x_u32[q_base + 2 * sub + 0] = src_qs[2 * sub + 0]; + x_u32[q_base + 2 * sub + 1] = src_qs[2 * sub + 1]; } - x_u32[i * MMQ_MMA_TILE_X_K_NVFP4 + 64 + kbx] = get_int_b1(bxi->d, 0); + x_u32[row_base + 64 + kbx] = get_int_b1(bxi->d, 0); } } @@ -1018,7 +997,7 @@ static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restric typedef tile<8, 8, int> tile_B; typedef tile<16, 8, float> tile_C; - constexpr int stride = mmq_get_mma_tile_x_k(type); + constexpr int stride = MMQ_MMA_TILE_X_K_FP4; constexpr int granularity = mmq_get_granularity_device(mmq_x); constexpr int rows_per_warp = 2 * granularity; constexpr int ntx = rows_per_warp / tile_C::I; @@ -1033,8 +1012,9 @@ static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restric // 2 threads per quad supply the packed scale register to the block_scale MMA, // see https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-block-scaling - const int tidx = threadIdx.x / 4 + (threadIdx.x % 2) * 8; - const int i0 = (threadIdx.y / ntx) * rows_per_warp; + const int tidx_A = threadIdx.x / 4 + (threadIdx.x % 2) * 8; + const int tidx_B = threadIdx.x / 4; + const int i0 = (threadIdx.y / ntx) * rows_per_warp; tile_A A[ntx][nfrags]; uint32_t scaleA[ntx][nfrags]; @@ -1045,7 +1025,7 @@ static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restric for (int frag = 0; frag < nfrags; ++frag) { const int k0 = k00 + frag * tile_A::J; load_ldmatrix(A[n][frag], x_qs + (i0 + n * tile_A::I) * stride + k0, stride); - scaleA[n][frag] = x_sc[(i0 + n * tile_A::I + tidx) * stride + k0 / tile_A::J]; + scaleA[n][frag] = x_sc[(i0 + n * tile_A::I + tidx_A) * stride + k0 / tile_A::J]; } } @@ -1058,7 +1038,7 @@ static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restric for (int frag = 0; frag < nfrags; ++frag) { const int k0 = frag * tile_B::J; load_generic(B[frag], y_qs + j0 * MMQ_TILE_Y_K + k0, MMQ_TILE_Y_K); - scaleB[frag] = y_sc[(j0 + tidx) * MMQ_TILE_Y_K + frag]; + scaleB[frag] = y_sc[(j0 + tidx_B) * MMQ_TILE_Y_K + frag]; } #pragma unroll @@ -1066,11 +1046,7 @@ static __device__ __forceinline__ void vec_dot_fp4_fp4_mma(const int * __restric #pragma unroll for (int frag = 0; frag < nfrags; ++frag) { tile_C C = {}; - if constexpr (type == GGML_TYPE_MXFP4) { - mma_block_scaled (C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); - } else { - mma_nvfp4_blackwell(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); - } + mma_block_scaled_fp4(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); #pragma unroll for (int l = 0; l < tile_C::ne; ++l) { sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; @@ -3487,8 +3463,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( #if defined(BLACKWELL_MMA_AVAILABLE) // FP4 tile stores 8 blocks - constexpr int ne_block = (type == GGML_TYPE_MXFP4) ? 8 * QK_MXFP4 : - (type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1; + constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1; #else constexpr int ne_block = 4 * QK8_1; #endif // defined(BLACKWELL_MMA_AVAILABLE) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 65d9298a407a..efe84f37f18a 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -77,9 +77,7 @@ static __global__ void quantize_mmq_nvfp4( const int64_t ne0, const int64_t ne1, const int64_t ne2) { #if defined(BLACKWELL_MMA_AVAILABLE) - const int lane_id = threadIdx.x & 31; - - const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * 8; + const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * QK_NVFP4_SUB; if (i0_base >= ne0) { return; } @@ -95,14 +93,16 @@ static __global__ void quantize_mmq_nvfp4( } const int64_t ib = blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x; - block_nvfp4_mmq * y = (block_nvfp4_mmq *) vy; - block_nvfp4_mmq * yb = y + ib; + block_fp4_mmq * y = (block_fp4_mmq *) vy; + block_fp4_mmq * yb = y + ib; + + const int sub = (i0_base % QK_K) / QK_NVFP4_SUB; - float vals_raw[8]; + float vals_raw[QK_NVFP4_SUB]; float amax_raw = 0.0f; const int64_t base_idx = i3 * s03 + i2 * s02 + i01 * s01; #pragma unroll - for (int k = 0; k < 8; k++) { + for (int k = 0; k < QK_NVFP4_SUB; k++) { const int64_t i00 = i0_base + k; if (i00 < ne00) { const float v = x[base_idx + i00]; @@ -113,31 +113,28 @@ static __global__ void quantize_mmq_nvfp4( } } - const float sub_max = fmaxf(amax_raw, __shfl_xor_sync(0xFFFFFFFFu, amax_raw, 1)); static constexpr int test_offsets[5] = { 0, -1, 1, -2, 2}; - const int first_fp8_code = (int) ggml_cuda_fp32_to_ue4m3(sub_max * (1.0f / 6.0f)); + const int first_fp8_code = (int) ggml_cuda_fp32_to_ue4m3(amax_raw * (1.0f / 6.0f)); float best_err = FLT_MAX; uint8_t fp8_code = 0; float subblock_scale = 0.0f; #pragma unroll // Check +/- 2 to find best code to reduce NVFP4 activation loss. Negligible overhead on Blackwell. - for (int i = 0; i < 5; i++) { + for (int i = 0; i < 1; i++) { const int test_code = first_fp8_code + test_offsets[i]; const uint8_t code = (uint8_t) test_code; const float test_scale = ggml_cuda_ue4m3_to_fp32(code); const float test_inv_scale = test_scale > 0.0f ? 0.5f / test_scale : 0.0f; float cur_err = 0.0f; #pragma unroll - for (int k = 0; k < 8; ++k) { + for (int k = 0; k < QK_NVFP4_SUB; ++k) { const float v = vals_raw[k]; const uint8_t q = ggml_cuda_float_to_fp4_e2m1(v, test_inv_scale); const float err_diff = fabsf(v) - fabsf(kvalues_mxfp4[q & 0x7]) * test_scale; cur_err = fmaf(err_diff, err_diff, cur_err); } - cur_err = cur_err + __shfl_xor_sync(__activemask(), cur_err, 1); - if (cur_err < best_err) { best_err = cur_err; fp8_code = test_code; @@ -146,16 +143,20 @@ static __global__ void quantize_mmq_nvfp4( } const float inv_scale = subblock_scale > 0.0f ? 0.5f / subblock_scale : 0.0f; - uint32_t q_word = 0; + uint32_t q0 = 0; + uint32_t q1 = 0; #pragma unroll // this is faster than the previous __nv_fp4x4_e2m1 - for (int k = 0; k < 8; ++k) { - q_word |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k], inv_scale) << (4 * k); + for (int k = 0; k < QK_NVFP4_SUB / 4; ++k) { + q0 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 0], inv_scale) << (8 * k); + q0 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 8], inv_scale) << (8 * k + 4); + q1 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 4], inv_scale) << (8 * k); + q1 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 12], inv_scale) << (8 * k + 4); } - yb->qs_u32[lane_id] = q_word; - if ((lane_id & 1) == 0) { - reinterpret_cast(yb->sc4_u32)[lane_id >> 1] = fp8_code; - } + uint32_t * yqs = reinterpret_cast(yb->qs); + yqs[2 * sub + 0] = q0; + yqs[2 * sub + 1] = q1; + reinterpret_cast(yb->d4)[sub] = fp8_code; #else NO_DEVICE_CODE; // This is for Blackwell NVFP4 activations only. #endif // defined(BLACKWELL_MMA_AVAILABLE) @@ -408,49 +409,37 @@ void quantize_mmq_q8_1_cuda( } } -void quantize_mmq_mxfp4_cuda(const float * x, - const int32_t * ids, - void * vy, - [[maybe_unused]] const ggml_type type_src0, - const int64_t ne00, - const int64_t s01, - const int64_t s02, - const int64_t s03, - const int64_t ne0, - const int64_t ne1, - const int64_t ne2, - const int64_t ne3, - cudaStream_t stream) { - GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0); - - constexpr int nwarps = 8; - constexpr int vals_per_warp = 2 * QK_MXFP4; - constexpr int vals_per_block = nwarps * vals_per_warp; - - const int64_t block_num_y = (ne0 + vals_per_block - 1) / vals_per_block; - const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); - const dim3 block_size(WARP_SIZE, nwarps, 1); - - quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); -} - -void quantize_mmq_nvfp4_cuda( +void quantize_mmq_fp4_cuda( const float * x, const int32_t * ids, void * vy, const ggml_type type_src0, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { - GGML_ASSERT(type_src0 == GGML_TYPE_NVFP4); + GGML_ASSERT(type_src0 == GGML_TYPE_MXFP4 || type_src0 == GGML_TYPE_NVFP4); GGML_ASSERT(ne00 % 8 == 0); GGML_ASSERT(ne0 > 0); - constexpr int nvfp4_block_size = 128; - const int64_t block_num_y = (ne0 + 8 * nvfp4_block_size - 1) / (8 * nvfp4_block_size); - const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); - const dim3 block_size(nvfp4_block_size, 1, 1); - if (ids) { - quantize_mmq_nvfp4<<>>( - x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + if (type_src0 == GGML_TYPE_NVFP4) { + constexpr int nvfp4_block_size = 128; + const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); + const dim3 block_size(nvfp4_block_size, 1, 1); + const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); + if (ids) { + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + } else { + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + } } else { - quantize_mmq_nvfp4<<>>( - x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); + GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0); + + constexpr int nwarps = 8; + constexpr int vals_per_warp = 2 * QK_MXFP4; + constexpr int vals_per_block = nwarps * vals_per_warp; + + const int64_t block_num_y = (ne0 + vals_per_block - 1) / vals_per_block; + const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); + const dim3 block_size(WARP_SIZE, nwarps, 1); + + quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); } } diff --git a/ggml/src/ggml-cuda/quantize.cuh b/ggml/src/ggml-cuda/quantize.cuh index 8558b45ceb3a..768a3ae6de6c 100644 --- a/ggml/src/ggml-cuda/quantize.cuh +++ b/ggml/src/ggml-cuda/quantize.cuh @@ -26,7 +26,7 @@ void quantize_mmq_q8_1_cuda( ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03, int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream); -void quantize_mmq_mxfp4_cuda(const float * x, +void quantize_mmq_fp4_cuda(const float * x, const int32_t * ids, void * vy, ggml_type type_src0, @@ -39,8 +39,3 @@ void quantize_mmq_mxfp4_cuda(const float * x, int64_t ne2, int64_t ne3, cudaStream_t stream); - -void quantize_mmq_nvfp4_cuda( - const float * x, const int32_t * ids, void * vy, - ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03, - int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream); From 78596bfa5496fe7b6be254450772616c7f458317 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 09:30:08 -0700 Subject: [PATCH 09/23] Updated block_fp4_mmq packing comment Co-authored-by: Aman Gupta --- ggml/src/ggml-cuda/mmq.cuh | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 405231d2951a..bc51fdbd0730 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -46,9 +46,12 @@ struct block_q8_1_mmq { int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each }; -struct block_fp4_mmq { // Used for both MXFP4[8 E8M0 scales (1 per 32)) and NVFP4 (16 UE4M3 scales (1 per 16)) - uint32_t d4[4]; // 2 packed per uint32: d4[0]={s0,s1}, d4[1]={s2,s3}, etc. - int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte), 8 blocks of 32 values +// this struct is used for fp4 data types (currently only used in Blackwell) +// mxfp4 has block size 32, each int32 of d4 contains 2 e8m0 scales in the lower 16 bits +// nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales +struct block_fp4_mmq { + uint32_t d4[4]; + int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte) }; static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); From a68327c7cb9dce7a26d07233ce20f5b80b4245df Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 09:36:35 -0700 Subject: [PATCH 10/23] Added assert for QK_K == 8 * QK_MXFP4 in mul_mat_q --- ggml/src/ggml-cuda/mmq.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 0a6f898394aa..17f5ce84941a 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -203,9 +203,10 @@ void ggml_cuda_mul_mat_q( quantize_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } - CUDA_CHECK(cudaGetLastError()); + CUDA_CHECK(cudaGetLastError()) ; } + GGML_ASSERT(QK_K == 8 * QK_MXFP4); const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; From 6e31a22bc2bf8b70455159a2abb21fff522382c0 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 09:37:56 -0700 Subject: [PATCH 11/23] Removed extra space typo --- ggml/src/ggml-cuda/mmq.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 17f5ce84941a..3b820172830f 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -203,7 +203,7 @@ void ggml_cuda_mul_mat_q( quantize_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } - CUDA_CHECK(cudaGetLastError()) ; + CUDA_CHECK(cudaGetLastError()); } GGML_ASSERT(QK_K == 8 * QK_MXFP4); From 58e277e44a0f5e313f8eab49f5ca51e1e76a1f08 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 10:02:00 -0700 Subject: [PATCH 12/23] Changed NVFP4 quant assert and using get_int_b4 --- ggml/src/ggml-cuda/mmq.cuh | 2 +- ggml/src/ggml-cuda/quantize.cu | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index bc51fdbd0730..bf12ba8b9a24 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -980,7 +980,7 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res x_u32[q_base + 2 * sub + 1] = src_qs[2 * sub + 1]; } - x_u32[row_base + 64 + kbx] = get_int_b1(bxi->d, 0); + x_u32[row_base + 64 + kbx] = get_int_b4(bxi->d, 0); } } diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index efe84f37f18a..06ebf052dd7f 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -121,7 +121,7 @@ static __global__ void quantize_mmq_nvfp4( float subblock_scale = 0.0f; #pragma unroll // Check +/- 2 to find best code to reduce NVFP4 activation loss. Negligible overhead on Blackwell. - for (int i = 0; i < 1; i++) { + for (int i = 0; i < 5; i++) { const int test_code = first_fp8_code + test_offsets[i]; const uint8_t code = (uint8_t) test_code; const float test_scale = ggml_cuda_ue4m3_to_fp32(code); @@ -414,10 +414,10 @@ void quantize_mmq_fp4_cuda( const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { GGML_ASSERT(type_src0 == GGML_TYPE_MXFP4 || type_src0 == GGML_TYPE_NVFP4); - GGML_ASSERT(ne00 % 8 == 0); GGML_ASSERT(ne0 > 0); if (type_src0 == GGML_TYPE_NVFP4) { + GGML_ASSERT(ne00 % QK_NVFP4 == 0); constexpr int nvfp4_block_size = 128; const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); const dim3 block_size(nvfp4_block_size, 1, 1); From 0e2c7948ed88d1d0a51bfef85f8130ae40fc12ee Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 10:08:28 -0700 Subject: [PATCH 13/23] Removed bool has_ids template from quantize --- ggml/src/ggml-cuda/quantize.cu | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 06ebf052dd7f..b3b6cb312c5b 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -70,7 +70,7 @@ __device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) { return static_cast(biased); } -template + static __global__ void quantize_mmq_nvfp4( const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, @@ -85,7 +85,7 @@ static __global__ void quantize_mmq_nvfp4( const int64_t i1 = blockIdx.x; const int64_t i2 = blockIdx.z % ne2; const int64_t i3 = blockIdx.z / ne2; - const int64_t i01 = has_ids ? ids[i1] : i1; + const int64_t i01 = ids ? ids[i1] : i1; const int64_t k_block = i0_base / QK_K; const int64_t blocks_per_col = (ne0 + QK_K - 1) / QK_K; if (k_block >= blocks_per_col) { @@ -422,13 +422,8 @@ void quantize_mmq_fp4_cuda( const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); const dim3 block_size(nvfp4_block_size, 1, 1); const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); - if (ids) { - quantize_mmq_nvfp4<<>>( - x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); - } else { - quantize_mmq_nvfp4<<>>( - x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); - } + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); } else { GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0); From 72fc01709aed655b707d544343fd5454da875345 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 10:14:10 -0700 Subject: [PATCH 14/23] Updated block_fp4_mmq packing comment Co-authored-by: Aman Gupta --- ggml/src/ggml-cuda/mmq.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 3b820172830f..89f307bdf6fb 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -147,7 +147,7 @@ void ggml_cuda_mul_mat_q( // Stride depends on quantization format const int64_t s12 = use_native_fp4 ? - ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : // block_fp4_mmq holds 256 values (8 blocks of 32) + ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : // block_fp4_mmq holds 256 values ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; From 7fcc8c076166372f4db2050d872c368bf189526f Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 10:21:17 -0700 Subject: [PATCH 15/23] Added ue4m3 bounds check for testscale --- ggml/src/ggml-cuda/quantize.cu | 3 +++ 1 file changed, 3 insertions(+) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index b3b6cb312c5b..d4b3b8758185 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -123,6 +123,9 @@ static __global__ void quantize_mmq_nvfp4( #pragma unroll // Check +/- 2 to find best code to reduce NVFP4 activation loss. Negligible overhead on Blackwell. for (int i = 0; i < 5; i++) { const int test_code = first_fp8_code + test_offsets[i]; + if (test_code < 0 || test_code > 0x7e) { + continue; + } const uint8_t code = (uint8_t) test_code; const float test_scale = ggml_cuda_ue4m3_to_fp32(code); const float test_inv_scale = test_scale > 0.0f ? 0.5f / test_scale : 0.0f; From 7c73198d28d650e528e633f106d9bdc510cbba98 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 12:55:34 -0700 Subject: [PATCH 16/23] Removed whitespace on line 52 of mmq.cuh --- ggml/src/ggml-cuda/mmq.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index bf12ba8b9a24..34aae42626b0 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -49,7 +49,7 @@ struct block_q8_1_mmq { // this struct is used for fp4 data types (currently only used in Blackwell) // mxfp4 has block size 32, each int32 of d4 contains 2 e8m0 scales in the lower 16 bits // nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales -struct block_fp4_mmq { +struct block_fp4_mmq { uint32_t d4[4]; int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte) }; From 6b26a1c7235bafa837b83eafad40310329b4b2bd Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Wed, 15 Apr 2026 21:48:19 -0700 Subject: [PATCH 17/23] Fixed MMQ_ITER_K_FP4 returning on non-FP4 models when running on Blackwell --- ggml/src/ggml-cuda/mmq.cuh | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 34aae42626b0..ede436fb3356 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -146,10 +146,11 @@ static int get_mmq_y_host(const int cc) { static constexpr __device__ int get_iter_k([[maybe_unused]] const ggml_type type) { #if defined(BLACKWELL_MMA_AVAILABLE) +if (type == GGML_TYPE_NVFP4 || type == GGML_TYPE_MXFP4) { return MMQ_ITER_K_FP4; -#else - return MMQ_ITER_K; +} #endif // defined(BLACKWELL_MMA_AVAILABLE) + return MMQ_ITER_K; } static constexpr __device__ int get_mmq_y_device() { From e34b6ff6d014073a7c01814156e103bb59fc5927 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Thu, 16 Apr 2026 10:55:47 -0700 Subject: [PATCH 18/23] Change GGML_ASSERT to static_assert Co-authored-by: Oliver Simons --- ggml/src/ggml-cuda/mmq.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 89f307bdf6fb..e1add5e03316 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -206,7 +206,7 @@ void ggml_cuda_mul_mat_q( CUDA_CHECK(cudaGetLastError()); } - GGML_ASSERT(QK_K == 8 * QK_MXFP4); + static_assert(QK_K == 8 * QK_MXFP4, "QK_K needs to be 8 * QK_MXFP4"); const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; From 02df26386eb46d9171d6e4bec1f25c723a4a4217 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Thu, 16 Apr 2026 11:01:26 -0700 Subject: [PATCH 19/23] Whitespace fixes --- ggml/src/ggml-cuda/mmq.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index ede436fb3356..420c77b5ca56 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -10,9 +10,9 @@ using namespace ggml_cuda_mma; #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. -#define MMQ_ITER_K 256 -#define MMQ_ITER_K_FP4 512 -#define MMQ_NWARPS 8 +#define MMQ_ITER_K 256 +#define MMQ_ITER_K_FP4 512 +#define MMQ_NWARPS 8 typedef void (*load_tiles_mmq_t)(const char * __restrict__ x, int * x_tile, const int kbx0, const int i_max, const int stride); typedef void (*vec_dot_mmq_t)(const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00); @@ -247,7 +247,7 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { #if defined(BLACKWELL_MMA_AVAILABLE) case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_FP4; #else - case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; + case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; #endif case GGML_TYPE_Q2_K: return MMQ_MMA_TILE_X_K_Q2_K; case GGML_TYPE_Q3_K: return MMQ_MMA_TILE_X_K_Q3_K; From 920459082ab6d9fd3d9cec52563c1155ffe4ca0a Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Thu, 16 Apr 2026 11:05:32 -0700 Subject: [PATCH 20/23] Change amax_raw mul 1/6 to: / 6 Co-authored-by: Oliver Simons --- ggml/src/ggml-cuda/quantize.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index d4b3b8758185..52f664719ae1 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -114,7 +114,7 @@ static __global__ void quantize_mmq_nvfp4( } static constexpr int test_offsets[5] = { 0, -1, 1, -2, 2}; - const int first_fp8_code = (int) ggml_cuda_fp32_to_ue4m3(amax_raw * (1.0f / 6.0f)); + const int first_fp8_code = (int) ggml_cuda_fp32_to_ue4m3(amax_raw / 6.0f); float best_err = FLT_MAX; uint8_t fp8_code = 0; From 667cc38d86352a76e2b9d305b9733fe5a65e3ddd Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Thu, 16 Apr 2026 11:13:43 -0700 Subject: [PATCH 21/23] Hoisted kbx0 and kbx out of the loop --- ggml/src/ggml-cuda/mmq.cuh | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 420c77b5ca56..6a80196d3e31 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -961,6 +961,9 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res const int kbx = txi % threads_per_row; const int row_in_warp = txi / threads_per_row; + const block_nvfp4 * bxi_base = (const block_nvfp4 *) x + kbx0 + kbx; + uint32_t * x_u32_scale = x_u32 + 64 + kbx; + #pragma unroll for (int i0 = 0; i0 < mmq_y; i0 += rows_per_warp * nwarps) { int i = i0 + threadIdx.y * rows_per_warp + row_in_warp; @@ -969,7 +972,7 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res i = min(i, i_max); } - const block_nvfp4 * bxi = (const block_nvfp4 *) x + kbx0 + i * stride + kbx; + const block_nvfp4 * bxi = bxi_base + i * stride; const int row_base = i * MMQ_MMA_TILE_X_K_FP4; const int q_base = row_base + 8 * kbx; @@ -981,7 +984,7 @@ static __device__ __forceinline__ void load_tiles_nvfp4_nvfp4(const char * __res x_u32[q_base + 2 * sub + 1] = src_qs[2 * sub + 1]; } - x_u32[row_base + 64 + kbx] = get_int_b4(bxi->d, 0); + x_u32_scale[row_base] = get_int_b4(bxi->d, 0); } } From 553c3a85c2e025f30b7ffcdd442d7f2f24047088 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Fri, 17 Apr 2026 13:39:15 -0700 Subject: [PATCH 22/23] Update ggml/src/ggml-cuda/mmq.cuh MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/mmq.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 6a80196d3e31..6501e1389506 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -46,7 +46,7 @@ struct block_q8_1_mmq { int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each }; -// this struct is used for fp4 data types (currently only used in Blackwell) +// this struct is used for fp4 data types (currently only used for Blackwell) // mxfp4 has block size 32, each int32 of d4 contains 2 e8m0 scales in the lower 16 bits // nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales struct block_fp4_mmq { From 0d9e0458ec0a3ac780a56dab7ef2ebd5a1f63c41 Mon Sep 17 00:00:00 2001 From: Michael Wand Date: Fri, 17 Apr 2026 14:19:56 -0700 Subject: [PATCH 23/23] Add endif blackwell mma comment MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/mmq.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 6501e1389506..3315fce5ce6d 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -248,7 +248,7 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_FP4; #else case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; -#endif +#endif // defined(BLACKWELL_MMA_AVAILABLE) case GGML_TYPE_Q2_K: return MMQ_MMA_TILE_X_K_Q2_K; case GGML_TYPE_Q3_K: return MMQ_MMA_TILE_X_K_Q3_K; case GGML_TYPE_Q4_K: return MMQ_MMA_TILE_X_K_Q8_1;