perf(cohere2): fuse the parallel-block residual adds with the next norm

[inureyes]

Summary

A Cohere2 block ends in (attn + mlp) + x and the next block starts with a LayerNorm of that sum, so on the decode critical path every layer boundary is three dependent kernels: two adds and MLX's layer_norm_single_row. Two commits remove two of them, both byte-identical to the current output:

• compiled_add3 runs the two adds as one compiled kernel with the same association order.
• fused_add3_layer_norm is one Metal kernel that writes the residual and its LayerNorm. It copies the pinned MLX single-row kernel (threadgroup size, 8 reads per thread, two-stage simd and threadgroup reductions, precise::rsqrt, the affine step with the bias read from memory as fast::layer_norm passes it), so both outputs match the unfused pair bit for bit. layers::residual_add3_layer_norm falls back to the pair off Metal, above a 6656-wide row, or on mixed dtypes, and MLXCEL_FUSED_ADD_NORM=0 forces the fallback. The Cohere2 layer loop fuses block i's add with block i+1's input norm, and the last block's add with the final norm.

Measurements

command-r7b 4-bit on M1 Ultra, 500-token prompt, 128 tokens, eight interleaved ABBA pairs, measured with #1947 applied:

Change	Off	On
compiled_add3 (budget pinned on both arms)	109.35 tok/s	110.54 tok/s (+1.1%)
fused add + LayerNorm (same binary, MLXCEL_FUSED_ADD_NORM=0 vs unset)	112.45	113.72 (+1.1%)

Prefill is unchanged.

Validation

• A unit test asserts exact equality with compiled_add3 followed by LayerNorm::forward across f16 and bf16, with and without a bias, several rows, and a width that is not a multiple of 8. Multiplying the kernel's rsqrt by 1.001 makes it fail, which shows the fused path runs and the test discriminates.
• Greedy generations for three prompts at 200 tokens are identical to main. The logit-trace gate was not run because the change is bitwise identical by construction and by that test.
• On the branch this was developed on, the full workspace gate (make verify-test, 11,412 passed), clippy and fmt pass; on this branch, clippy, fmt and the new test.

Two mx.fast.metal_kernel details worth knowing: a 0-d input is passed to the kernel as a scalar rather than a pointer, and a small input can land in the constant address space, so the kernel takes the missing bias as a one-element array and uses auto pointers. Upstream's kernel starts one tail loop at a negative index for threads past the end of a narrow row; the copy clamps the start to 0, which leaves the result unchanged.