magetypes: mul_add is fused on v3/neon but unfused on scalar tier (and wasm128) — cross-tier output divergence

@lilith is already working on this.

Since Jul 17, 2026.

The generic `mul_add` lowers to `_mm256_fmadd_ps`/`_mm_fmadd_ps` on x86 v3/v4 and vfma on NEON (exactly fused), but the scalar tier maps to `nostd_math::fmaf`, which is a documented non-fused fallback (`a * b + c`), and wasm128 has no FMA either. A `#[magetypes]` kernel written once therefore computes different results per tier on inputs where fused-vs-unfused rounding straddles a boundary.

Real-world hit: imazen/zenavif's float YUV kernels — the i686 CI leg (which exercises the scalar tier) failed a byte-identity test against a scalar reference on a 1-LSB green-channel mismatch (1 pixel in 3072; red since 2026-07-06). Fixed on the zenavif side by removing `mul_add` from determinism-sensitive kernels (zenavif@96200c2e).

Not proposing a behavior change — hardware without FMA can't fuse, and software-exact fmaf is slow — but two things would help consumers:

1. Document on `mul_add` (and `mul_sub`) that fuse-ness is tier-dependent, so output-byte-determinism kernels must use unfused `a * b + c` (which is exact per-op IEEE and identical on every tier). `round()` is already consistent (ties-even everywhere, including scalar `roundevenf` — that consistency note on scalar `round()` is a good model).
2. Optionally: a doc note that the i686/scalar path is the determinism canary — nothing else in CI typically exercises the scalar tier on x86.

For reference, the scalar fallback: `nostd_math::fmaf(a, b, c) = a * b + c` ("non-fused fallback").