Metal: `QQMatmul::eval_gpu` gemv path silently drops `global_scale_x`/`global_scale_w` for nvfp4

[Brooooooklyn]

Summary

QQMatmul::eval_gpu on Metal silently drops the global_scale_x / global_scale_w inputs when running qqmm in nvfp4 mode along the gemv special case (x.shape(-2) == 1 with pre-quantized w). The general case correctly throws "[QQMatmul] NYI for the general case", but the gemv branch runs to completion using only per-group scales — producing numerically incorrect results when tensor-scale nvfp4 weights are actually in use.

The sibling APIs quantize() and dequantize() already guard this on Metal:

// mlx/ops.cpp:4940-4945  (quantize)
if (to_stream(s).device == Device::gpu && metal::is_available() &&
    global_scale.has_value()) {
  throw std::invalid_argument("[quantize] Global scale is not supported on the Metal backend.");
}
// mlx/ops.cpp:5205-5210 has the same guard for dequantize().

…but qqmm() has no equivalent guard, so the global scales get packed into inputs (ops.cpp:4666-4669) and propagate to Metal where the gemv path ignores them.

Where the inputs get dropped

mlx/backend/metal/quantized.cpp:1611-1655:

void QQMatmul::eval_gpu(const std::vector<array>& inputs, array& out) {
  // ...
  bool w_quantized = (inputs[1].dtype() == uint32);
  if (w_quantized && inputs[0].shape(-2) == 1) {
    // ...
    quantize_dequantize(x, xhat, mode, group_size_, bits_, d, s); // <-- no global_scale_x
    // ...
    dispatch_qmv(xhat, w, scales, std::nullopt, out, group_size_, bits_, M, N, K, d, s, mode);
    //                                ^^^^^^^^^^                                        ^^^^
    //                            biases slot                              no global_scale_w
    return;
  } else {
    throw std::runtime_error("[QQMatmul] NYI for the general case");
  }
}

inputs[3] (global_scale_x) and inputs[4] (global_scale_w) — packed by ops.cpp when both are provided for nvfp4 — are never read on this branch.

For comparison, the CUDA backend reads them in mlx/backend/cuda/quantized/qqmm.cpp:93,134-135:

// CUDA reads from the tail of inputs[]
global_scale = inputs[inputs.size() - 2];                   // qqmm.cpp:93
global_scale_x = inputs[inputs.size() - 2];                 // qqmm.cpp:134
global_scale_w = inputs[inputs.size() - 1];                 // qqmm.cpp:135

Impact

Tensor-scale nvfp4 (introduced in #3022) is intentionally CPU/CUDA-only — That decision was carried out for quantize / dequantize via the explicit Metal guards above, but qqmm was missed: the gemv path on Metal does not throw, it just silently produces wrong outputs.

Suggested fix

Mirror the general case's behavior — throw NYI from the gemv path when global scales are present — until proper tensor-scale-aware Metal kernels land. Something like:

void QQMatmul::eval_gpu(const std::vector<array>& inputs, array& out) {
  // Reject tensor-scale nvfp4 on Metal until the kernels support it.
  if (mode_ == QuantizationMode::Nvfp4 && inputs.size() > 3) {
    throw std::runtime_error(
        "[QQMatmul] nvfp4 with tensor-scale (global_scale_x/w) is not supported on the Metal backend.");
  }
  // ...existing body...
}

Alternatively, add the API-level guard in qqmm() in ops.cpp next to the existing quantize / dequantize ones so the error is raised at graph-construction time rather than during evaluation.

Context

Found while bringing up tensor-scale nvfp4 in a downstream project. We work around it by sticking to per-group nvfp4 (no global_scale) for Metal-only inference paths, which is unaffected. Filing this so the silent-incorrect behavior is either gated with an explicit error or eventually implemented in Metal.