Skip to content

metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels (mirror of #330 by @Lyxot) #348

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
danielhanchen wants to merge 2 commits into
base: base/upstream-ee17f293b
Choose a base branch
from
Open

metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels (mirror of #330 by @Lyxot) #348

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3adf93f
metal : add NVFP4 mul_mat, mul_mat_id and get_rows kernels
Lyxot Oct 27, 2026
56ec1e8
Fix merge conflicts
danielhanchen Nov 3, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions ggml/src/ggml-metal/ggml-metal-common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,6 +85,7 @@ static bool ggml_metal_mul_mv_mma_type_supported(enum ggml_type type) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q2_0:
case GGML_TYPE_MXFP4:
case GGML_TYPE_NVFP4:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_IQ2_XXS:
Expand Down
12 changes: 12 additions & 0 deletions ggml/src/ggml-metal/ggml-metal-device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1104,6 +1104,12 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
nsg = N_SG_MXFP4;
nr0 = N_R0_MXFP4;
} break;
case GGML_TYPE_NVFP4:
{
nsg = N_SG_NVFP4;
nr0 = N_R0_NVFP4;
smem = 32*sizeof(float);
} break;
case GGML_TYPE_Q2_K:
{
nsg = N_SG_Q2_K;
Expand Down Expand Up @@ -1399,6 +1405,12 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
nsg = N_SG_MXFP4;
nr0 = N_R0_MXFP4;
} break;
case GGML_TYPE_NVFP4:
{
nsg = N_SG_NVFP4;
nr0 = N_R0_NVFP4;
smem = 32*sizeof(float);
} break;
case GGML_TYPE_Q2_K:
{
nsg = N_SG_Q2_K;
Expand Down
4 changes: 1 addition & 3 deletions ggml/src/ggml-metal/ggml-metal-device.m
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1517,7 +1517,6 @@ static bool ggml_metal_supports_mul_mat_op(
bool mm_path) {
if (!has_simdgroup_reduction ||
op->src[0]->type == GGML_TYPE_F8_E4M3 ||
op->src[0]->type == GGML_TYPE_NVFP4 ||
op->src[0]->type == GGML_TYPE_TQ1_0) {
return false;
}
Expand Down Expand Up @@ -1966,8 +1965,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
};
}
case GGML_OP_GET_ROWS:
return op->src[0]->type != GGML_TYPE_NVFP4 &&
op->src[0]->type != GGML_TYPE_TQ1_0 &&
return op->src[0]->type != GGML_TYPE_TQ1_0 &&
op->src[0]->type != GGML_TYPE_F8_E4M3;
case GGML_OP_SET_ROWS:
{
Expand Down
3 changes: 3 additions & 0 deletions ggml/src/ggml-metal/ggml-metal-impl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,9 @@
#define N_R0_MXFP4 2
#define N_SG_MXFP4 2

#define N_R0_NVFP4 2
#define N_SG_NVFP4 2

#define N_R0_Q2_K 4
#define N_SG_Q2_K 2

Expand Down
1 change: 1 addition & 0 deletions ggml/src/ggml-metal/ggml-metal-ops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2765,6 +2765,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) {
op->src[0]->type == GGML_TYPE_Q5_1 ||
op->src[0]->type == GGML_TYPE_Q8_0 ||
op->src[0]->type == GGML_TYPE_MXFP4 ||
op->src[0]->type == GGML_TYPE_NVFP4 ||
op->src[0]->type == GGML_TYPE_IQ4_NL ||
false) && (ne11 >= 2 && ne11 <= 8)
) ||
Expand Down
8 changes: 8 additions & 0 deletions ggml/src/ggml-metal/kernels/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,6 +68,14 @@ static inline float e8m0_to_fp32(uint8_t x) {
return as_type<float>(bits);
}

// unsigned E4M3 (bias 7), 0x7F (NaN) decodes to 0 as on the CPU; unlike ggml_ue4m3_to_fp32 there is no 0.5 factor since kvalues_mxfp4_f is not doubled
// the 7 bits placed in an fp16 give the right value scaled by 2^-8, subnormals included
static inline float ue4m3_to_fp32(uint8_t x) {
const float v = (float) as_type<half>((ushort) ((x & 0x7F) << 7)) * 256.0f;

return x == 0x7F ? 0.0f : v;
}

static inline float dot(float x, float y) {
return x*y;
}
Expand Down
37 changes: 37 additions & 0 deletions ggml/src/ggml-metal/kernels/dequantize.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -378,6 +378,43 @@ void dequantize_mxfp4_t4(device const block_mxfp4 * xb, short il, thread type4 &
reg[3] = d * kvalues_mxfp4_f[(q2[4*il4 + 3] >> shr) & 0x0F];
}

// il selects one of the 4 sub-blocks; its 8 bytes hold elements 0..7 in the low nibbles and 8..15 in the high nibbles
template <typename type4x4>
void dequantize_nvfp4(device const block_nvfp4 * xb, short il, thread type4x4 & reg) {
device const uint8_t * q2 = xb->qs + il*(QK_NVFP4_SUB/2);

const float d = ue4m3_to_fp32(xb->d[il]);

float4x4 reg_f;

for (int i = 0; i < 4; ++i) {
const uint8_t shr = i >= 2 ? 4 : 0;

reg_f[i][0] = d * kvalues_mxfp4_f[(q2[4*(i%2) + 0] >> shr) & 0x0F];
reg_f[i][1] = d * kvalues_mxfp4_f[(q2[4*(i%2) + 1] >> shr) & 0x0F];
reg_f[i][2] = d * kvalues_mxfp4_f[(q2[4*(i%2) + 2] >> shr) & 0x0F];
reg_f[i][3] = d * kvalues_mxfp4_f[(q2[4*(i%2) + 3] >> shr) & 0x0F];
}

reg = (type4x4) reg_f;
}

template <typename type4>
void dequantize_nvfp4_t4(device const block_nvfp4 * xb, short il, thread type4 & reg) {
const short sub = il/4;
const short i = il%4;

device const uint8_t * q2 = xb->qs + sub*(QK_NVFP4_SUB/2) + 4*(i%2);

const float d = ue4m3_to_fp32(xb->d[sub]);
const uint8_t shr = i >= 2 ? 4 : 0;

reg[0] = d * kvalues_mxfp4_f[(q2[0] >> shr) & 0x0F];
reg[1] = d * kvalues_mxfp4_f[(q2[1] >> shr) & 0x0F];
reg[2] = d * kvalues_mxfp4_f[(q2[2] >> shr) & 0x0F];
reg[3] = d * kvalues_mxfp4_f[(q2[3] >> shr) & 0x0F];
}

template <typename type4x4>
void dequantize_q2_K(device const block_q2_K *xb, short il, thread type4x4 & reg) {
const float d = xb->d;
Expand Down
4 changes: 4 additions & 0 deletions ggml/src/ggml-metal/kernels/mul_mm.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -905,6 +905,7 @@ template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_m
template [[host_name("kernel_mul_mm_q5_1_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_q8_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_mxfp4_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_nvfp4_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_nvfp4, 4, dequantize_nvfp4, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_q2_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_q3_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_q4_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, float, float2x4>;
Expand All @@ -931,6 +932,7 @@ template [[host_name("kernel_mul_mm_q5_0_f16")]] kernel mul_mm_t kernel_mul_m
template [[host_name("kernel_mul_mm_q5_1_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_q8_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_mxfp4_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_nvfp4_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_nvfp4, 4, dequantize_nvfp4, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_q2_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_q3_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_q4_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, half, half2x4>;
Expand Down Expand Up @@ -970,6 +972,7 @@ template [[host_name("kernel_mul_mm_id_mxfp4_f32")]] kernel mul_mm_id kernel_m
#else
template [[host_name("kernel_mul_mm_id_mxfp4_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, float, float2x4>;
#endif
template [[host_name("kernel_mul_mm_id_nvfp4_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_nvfp4, 4, dequantize_nvfp4, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_id_q2_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_id_q3_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_id_q4_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, float, float2x4>;
Expand Down Expand Up @@ -1000,6 +1003,7 @@ template [[host_name("kernel_mul_mm_id_mxfp4_f16")]] kernel mul_mm_id kernel_m
#else
template [[host_name("kernel_mul_mm_id_mxfp4_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, half, half2x4>;
#endif
template [[host_name("kernel_mul_mm_id_nvfp4_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_nvfp4, 4, dequantize_nvfp4, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_id_q2_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_id_q3_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, half, half2x4>;
template [[host_name("kernel_mul_mm_id_q4_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, half, half2x4>;
Expand Down
105 changes: 105 additions & 0 deletions ggml/src/ggml-metal/kernels/mul_mv.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -894,6 +894,11 @@ template [[host_name("kernel_mul_mv_ext_mxfp4_f32_r1_3")]] kernel mul_mv_ext_q4
template [[host_name("kernel_mul_mv_ext_mxfp4_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_mxfp4, 32, dequantize_mxfp4_t4>;
template [[host_name("kernel_mul_mv_ext_mxfp4_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_mxfp4, 32, dequantize_mxfp4_t4>;

template [[host_name("kernel_mul_mv_ext_nvfp4_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_nvfp4, 64, dequantize_nvfp4_t4>;
template [[host_name("kernel_mul_mv_ext_nvfp4_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_nvfp4, 64, dequantize_nvfp4_t4>;
template [[host_name("kernel_mul_mv_ext_nvfp4_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_nvfp4, 64, dequantize_nvfp4_t4>;
template [[host_name("kernel_mul_mv_ext_nvfp4_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_nvfp4, 64, dequantize_nvfp4_t4>;

template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
Expand Down Expand Up @@ -3146,6 +3151,105 @@ kernel void kernel_mul_mv_mxfp4_f32(
kernel_mul_mv_mxfp4_f32_impl<N_R0_MXFP4, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
}

template<int NR0, typename args_t>
void kernel_mul_mv_nvfp4_f32_impl(
args_t args,
device const char * src0,
device const char * src1,
device char * dst,
threadgroup char * shmem,
uint3 tgpig,
ushort tiisg,
ushort sgitg) {
const short NSG = FC_mul_mv_nsg;

threadgroup float * shmem_f32 = (threadgroup float *) shmem;

const int r0 = tgpig.x;
const int r1 = tgpig.y;
const int im = tgpig.z;

const int first_row = (r0 * NSG + sgitg) * NR0;

const uint i12 = im%FC_mul_mv_ne12;
const uint i13 = im/FC_mul_mv_ne12;

const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;

device const block_nvfp4 * x = (device const block_nvfp4 *) (src0 + offset0);
device const float * y = (device const float *) (src1 + offset1);

const int nb = args.ne00/QK_NVFP4;
const int ns01 = args.nb01/args.nb00; // this can be larger than nb for permuted src0 tensors

const short ix = tiisg/2; // 0...15
const short it = tiisg%2; // 0 or 1, each thread takes 2 of the 4 sub-blocks

shmem_f32[tiisg] = kvalues_mxfp4_f[tiisg%16];
threadgroup_barrier(mem_flags::mem_threadgroup);

float4 yl[8];
float sumf[NR0]={0.f};

device const float * yb = y + ix*QK_NVFP4 + it*2*QK_NVFP4_SUB;

for (int ib = ix; ib < nb && ib < ns01; ib += 16) {
device const float4 * y4 = (device const float4 *) yb;

FOR_UNROLL (short i = 0; i < 8; ++i) {
yl[i] = y4[i];
}

FOR_UNROLL (short row = 0; row < NR0; row++) {
device const block_nvfp4 & xb = x[row*ns01 + ib];
device const uint8_t * q2 = xb.qs + 2*it*(QK_NVFP4_SUB/2);

float sumq = 0.f;

FOR_UNROLL (short s = 0; s < 2; ++s) {
device const uint8_t * q = q2 + 8*s;

float4 acc1 = yl[4*s + 0]*float4(shmem_f32[q[0] & 0x0F], shmem_f32[q[1] & 0x0F], shmem_f32[q[2] & 0x0F], shmem_f32[q[3] & 0x0F]);
float4 acc2 = yl[4*s + 1]*float4(shmem_f32[q[4] & 0x0F], shmem_f32[q[5] & 0x0F], shmem_f32[q[6] & 0x0F], shmem_f32[q[7] & 0x0F]);
float4 acc3 = yl[4*s + 2]*float4(shmem_f32[q[0] >> 4 ], shmem_f32[q[1] >> 4 ], shmem_f32[q[2] >> 4 ], shmem_f32[q[3] >> 4 ]);
float4 acc4 = yl[4*s + 3]*float4(shmem_f32[q[4] >> 4 ], shmem_f32[q[5] >> 4 ], shmem_f32[q[6] >> 4 ], shmem_f32[q[7] >> 4 ]);

acc1 = (acc1 + acc3) + (acc2 + acc4);

sumq += ue4m3_to_fp32(xb.d[2*it + s]) * ((acc1[0] + acc1[1]) + (acc1[2] + acc1[3]));
}

sumf[row] += sumq;
}

yb += 16 * QK_NVFP4;
}

device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;

for (int row = 0; row < NR0 && first_row + row < args.ne0; ++row) {
float sum_all = simd_sum(sumf[row]);
if (tiisg == 0) {
dst_f32[first_row + row] = sum_all;
}
}
}

[[host_name("kernel_mul_mv_nvfp4_f32")]]
kernel void kernel_mul_mv_nvfp4_f32(
constant ggml_metal_kargs_mul_mv & args,
device const char * src0,
device const char * src1,
device char * dst,
threadgroup char * shmem [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]]) {

kernel_mul_mv_nvfp4_f32_impl<N_R0_NVFP4, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
}

template<int nr0, typename args_t>
void kernel_mul_mv_tq2_0_f32_impl(
args_t args,
Expand Down Expand Up @@ -3412,6 +3516,7 @@ template [[host_name("kernel_mul_mv_id_q5_0_f32")]] kernel kernel_mul_mv_id_t
template [[host_name("kernel_mul_mv_id_q5_1_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<mul_vec_q_n_f32_impl<block_q5_1, N_R0_Q5_1>>>;

template [[host_name("kernel_mul_mv_id_mxfp4_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_mxfp4_f32_impl<N_R0_MXFP4>>>;
template [[host_name("kernel_mul_mv_id_nvfp4_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_nvfp4_f32_impl<N_R0_NVFP4>>>;

template [[host_name("kernel_mul_mv_id_q2_K_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_q2_K_f32_impl <N_R0_Q2_K>>>;
template [[host_name("kernel_mul_mv_id_q3_K_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_q3_K_f32_impl <N_R0_Q3_K>>>;
Expand Down
1 change: 1 addition & 0 deletions ggml/src/ggml-metal/kernels/mul_mv_mma.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -611,6 +611,7 @@ MUL_MV_MMA_GEN("bf16", bfloat4x4, 1, dequantize_bf16)
MUL_MV_MMA_GEN("q1_0", block_q1_0, 8, dequantize_q1_0)
MUL_MV_MMA_GEN("q2_0", block_q2_0, 4, dequantize_q2_0)
MUL_MV_MMA_GEN("mxfp4", block_mxfp4, 2, dequantize_mxfp4)
MUL_MV_MMA_GEN("nvfp4", block_nvfp4, 4, dequantize_nvfp4)
MUL_MV_MMA_GEN("q2_K", block_q2_K, QK_NL, dequantize_q2_K)
MUL_MV_MMA_GEN("q3_K", block_q3_K, QK_NL, dequantize_q3_K)
MUL_MV_MMA_GEN("iq2_xxs", block_iq2_xxs, QK_NL, dequantize_iq2_xxs)
Expand Down
1 change: 1 addition & 0 deletions ggml/src/ggml-metal/kernels/quantize.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -341,6 +341,7 @@ template [[host_name("kernel_get_rows_q5_0")]] kernel get_rows_q_t kernel_get
template [[host_name("kernel_get_rows_q5_1")]] kernel get_rows_q_t kernel_get_rows_q<block_q5_1, 2, dequantize_q5_1>;
template [[host_name("kernel_get_rows_q8_0")]] kernel get_rows_q_t kernel_get_rows_q<block_q8_0, 2, dequantize_q8_0>;
template [[host_name("kernel_get_rows_mxfp4")]] kernel get_rows_q_t kernel_get_rows_q<block_mxfp4, 2, dequantize_mxfp4>;
template [[host_name("kernel_get_rows_nvfp4")]] kernel get_rows_q_t kernel_get_rows_q<block_nvfp4, 4, dequantize_nvfp4>;
template [[host_name("kernel_get_rows_q2_K")]] kernel get_rows_q_t kernel_get_rows_q<block_q2_K, QK_NL, dequantize_q2_K>;
template [[host_name("kernel_get_rows_q3_K")]] kernel get_rows_q_t kernel_get_rows_q<block_q3_K, QK_NL, dequantize_q3_K>;
template [[host_name("kernel_get_rows_q4_K")]] kernel get_rows_q_t kernel_get_rows_q<block_q4_K, QK_NL, dequantize_q4_K>;
Expand Down