[danielhanchen]

Mirror of #330 by @Lyxot, so the change can be kept mergeable from this repo. All credit for the kernels goes to @Lyxot.

• 3adf93fa5 is the original commit from metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels #330, unchanged, with its original author (Lyxot).
• 56ec1e88a merges b12065 into it. The only conflict was upstream Quant: OCP FP8 E4M3 support ggml-org/llama.cpp#27512 (OCP FP8 E4M3), which added GGML_TYPE_F8_E4M3 to the same two Metal "unsupported" lists this PR removes GGML_TYPE_NVFP4 from. The resolution keeps upstream's F8_E4M3 entries and drops only NVFP4, so against b12065 the diff is exactly metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels #330's change (11 files, +174/-3, same as metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels #330 against its own base).

The nightly pin moves from #330 to this PR at 56ec1e88a; #330 stays open.

---

Original description from #330:

Fixes the slow generation in unslothai/unsloth#15409 (MiniMax-M3.1 GGUFs on Apple Silicon run at ~8.5 tok/s with "61/61 layers offloaded to GPU").

Problem

The MiniMax-M3.1 exports keep the routed experts in native NVFP4 (171 tensors, ~218 GiB in UD-Q8_K_XL, also present in UD-Q4_K_XL). ggml_metal_supports_mul_mat_op returns false for GGML_TYPE_NVFP4, and GET_ROWS rejects it too, so the scheduler leaves every expert MUL_MAT_ID on the CPU. The layer count still reports full offload because the tensors' layers are assigned to Metal; only the expert ops fall back. The reporter's log shows 218.71 GiB CPU_Mapped against 22.57 GiB MTL0_Mapped.

Upstream has no Metal NVFP4 kernels yet. ggml-org#20456 adds them but is open and written against the old single ggml-metal.metal, so it no longer applies.

Fix

Metal kernels for NVFP4 in every path MUL_MAT / MUL_MAT_ID / GET_ROWS can take, following the MXFP4 wiring:

• ue4m3_to_fp32 in kernels/common.h: places the 7 scale bits into an fp16 and rescales by 2^8, which covers UE4M3 subnormals exactly; 0x7F decodes to 0 as on the CPU. No 0.5 factor, because Metal's kvalues_mxfp4_f are the real E2M1 values while the CPU table is doubled.
• dequantize_nvfp4 (16 weights per call, nl = 4) and dequantize_nvfp4_t4 in kernels/dequantize.h, matching the CPU layout: sub-block s uses qs[8s..8s+7], low nibble j -> element j, high nibble -> element j+8.
• kernel_mul_mv_nvfp4_f32 (each lane pair covers one 64-weight block, each lane two contiguous sub-blocks) plus its mul_mv_id instantiation.
• mul_mm, mul_mm_id, get_rows, mul_mv_ext instantiations, and registration in the few-row MMA kernels (mul_mv_mma, ggml_metal_mul_mv_mma_type_supported). Without MMA, 2..8-row batches went to the ext kernel at ~3x the MXFP4 time.
• supports_op no longer rejects NVFP4 for MUL_MAT, MUL_MAT_ID and GET_ROWS. TQ1_0 is still rejected.

Evidence

M3 Max, base d854d98 (no Metal source differs from the b11887 sync point).

test-backend-ops test -b MTL0 -o MUL_MAT,MUL_MAT_ID,GET_ROWS -p nvfp4
126/126 tests passed   (base: all NVFP4 cases "not supported" on MTL0)

The cases cover mat-vec (n=1), MMA (n=2..16), mat-mat (n=16, 64), mul_mv_id, mul_mm_id and get_rows. Breaking the per-sub-block scale index, the nibble order, or the ext byte group each fails between 15 and 30 cases. The scale decoder was compared against the CPU formula for all 256 byte values (0 mismatches). That check was repeated on the compiled Metal kernel, because random test data never produces subnormal or 0x7F scales.

test-backend-ops perf, MUL_MAT m=4096 k=14336, src1 f32:

n	NVFP4 CPU	NVFP4 Metal	MXFP4 Metal
1	267 us	100 us	92 us
8	1925 us	152 us	159 us
512	118 ms	5.9 ms	5.9 ms

NVFP4 is 4.5 bits/weight against 4.25 for MXFP4, so the n=1 gap is roughly the extra bytes.

End to end

I took Qwen3-0.6B (BF16) and quantized all 196 attention/MLP projections to NVFP4 in ModelOpt layout: per-tensor scale_2 = amax/(6*448), an FP8 E4M3 scale per 16 values, round-to-nearest E2M1. The checkpoint was converted with this tree's convert_hf_to_gguf.py, which wrote all 196 tensors as NVFP4. The "before" row restores the old NVFP4 rejection in supports_op on the same build.

llama-perplexity, wikitext-2 test, -c 512 --chunks 40, M3 Max:

run	NVFP4 weights in	PPL	s/pass
BF16, Metal	-	19.4502	0.28
NVFP4, -ngl 99, before (log says "offloaded 29/29 layers")	CPU_Mapped 532.79 MiB	22.4646	1.57
NVFP4, -ngl 99, this PR	MTL0_Mapped 533.60 MiB	22.4644	0.41
NVFP4, --device none	CPU_Mapped	22.4650	2.99

The "before" row has the same signature as the issue: every layer reported offloaded, but the NVFP4 weights stay on the CPU. With this PR, the only CPU_Mapped buffer left is the BF16 token embedding (296.75 MiB). Greedy generation: 187 tok/s on Metal vs 66 tok/s CPU only. The NVFP4-vs-BF16 PPL gap comes from the naive round-to-nearest quantizer used to make the checkpoint; the three NVFP4 runs agree to 4 significant figures.

Not measured on MiniMax-M3.1 itself (259 GB). Its expert ops are the MUL_MAT_ID / GET_ROWS shapes covered by the op tests above.

Notes

The ext kernel (dequantize_nvfp4_t4) is not reached by test-backend-ops once MMA handles 2..8 rows. It still runs when MMA is ineligible (no simdgroup_mm, src1 strides not 16-byte aligned), the same as MXFP4. It passed the full suite before MMA was registered.

Pinning into the nightly needs a pr-set.json entry for this PR's commit.

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-11-04T11:27:02.756649Z	56ec1e8	PR opened
🔒 Security Review	✅ Completed 2026-11-04T11:28:22.107114Z	56ec1e8	PR opened

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.