Skip to content

metal: add NVFP4 mul_mat, mul_mat_id and get_rows kernels - #330

Open
Lyxot wants to merge 1 commit into
unslothai:base/upstream-d854d98bafrom
Lyxot:fix/metal-nvfp4
Open

Lyxot wants to merge 1 commit into
unslothai:base/upstream-d854d98bafrom
Lyxot:fix/metal-nvfp4

Conversation

@Lyxot

@Lyxot Lyxot commented Oct 27, 2026 •

Copy link
Copy Markdown

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.

@Lyxot

Lyxot commented Oct 27, 2026

Copy link
Copy Markdown
Author

@codex review

@chatgpt-codex-connector

chatgpt-codex-connector Bot commented Oct 27, 2026 •

Copy link
Copy Markdown

Codex Review Summary

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

Review Status Commit Review trigger
📝 Code Review ⚠️ Failed 2026-10-28T06:08:19.498664Z
We couldn’t complete this review. Comment @codex review to try again.
3adf93f Manual request
ℹ️ 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.

@chatgpt-codex-connector

Copy link
Copy Markdown

Codex Review: Something went wrong. Try again later by commenting “@codex review”.

An unknown error occurred
ℹ️ 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".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

@Lyxot

Lyxot commented Oct 28, 2026

Copy link
Copy Markdown
Author

@codex review

@chatgpt-codex-connector

Copy link
Copy Markdown

Codex Review: Something went wrong. Try again later by commenting “@codex review”.

An unknown error occurred
ℹ️ 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".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant