Would new fused ops be welcome in mx.fast?

Every Llama-family decoder block does this twice per layer:

h = x + residual
y = mx.fast.rms_norm(h, w, eps)

That is two passes over the hidden state. mx.fast.rms_norm is already a tuned kernel, but
added ahead of it materialises h and reads it straight back.
fusing them into one pass reads x and residual, writes h and y, is done.
vLLM ships fused_add_rms_norm for this reason.

I wrote one via mx.fast.metal_kernel to find out whether it is worth anything on Apple silicon. Numbers below are fp16 on an M3 MacBook Air (81 GB/s measured copy bandwidth), median of 40 runs, compared against the current two-call path rather than against naive ops:

Output matches the two-call path to 7.8e-3 in fp16, which is one rounding of the final store.

Being straight about the shape of this:
it is a prefill-only win.
At a single row it is slower, because at that size the dispatch dominates and there is no bandwidth to save.

So it would want either a row-count threshold or documentation that it is for batched shapes.

Is this something you would want in mx.fast?
I have a working kernel and am happy to open a PR with tests if so, but
I did not want to send API surface you had not asked for. Equally happy to hear that ~1.2x on one op is below the bar.

Measurement setup: https://github.com/Rahu378/hardware-aware-gateway

Start with the existing mx.fast.rms_norm entry point and the mx.fast.metal_kernel approach described in the issue. Reproduce the reported batched and single-row benchmarks, then evaluate the API and whether a row-count threshold or documentation is needed. Done means a reviewed fused operation with tests confirming output agreement and its performance tradeoffs.