Module Ir.Numerics

Allow tensor-core matmuls over uniform-f32 operands to compute in tf32 on backends with a tf32 tile shape (CUDA sm_80+): f32's exponent range with a 10-bit mantissa, accumulation in f32. Off by default — opt-in like PyTorch's allow_tf32, because enabling it silently changes numerics. Metal (simdgroup_float8x8 is genuine f32) and HIP (RDNA WMMA has no tf32-like shape) are unaffected.

Run the arithmetic over narrow-float storage (bf16, fp16, fp8) in f32 on backends that have no native narrow arithmetic — the CPU backends, where every narrow operator is an explicit widen/op/narrow round-trip anyway (gh-ocannl-517). Storage stays narrow: reads widen once at the load, the result narrows once at the store, and the intermediates of an assignment keep f32 mantissa instead of being rounded per operator. That is both the faithful reading of "16-bit storage with f32 compute" and what makes the vectorized renderings — which are f32/f64 shaped — reachable for narrow-storage kernels.

On by default: it strictly increases accuracy relative to per-operator rounding and is the precondition for narrow storage being a speedup rather than a pessimization on CPU. Turn it off to recover the pre-gh-517 semantics, where every operator rounds to the target node's storage precision. The GPU backends are unaffected either way — they have native 16-bit types and arithmetic, so their compute precision is their storage precision.

Compute fp16 in fp16 on CPU targets that have native 16-bit arithmetic (ARMv8.2-FP16, AVX512-FP16), instead of widening to f32 (gh-ocannl-516). This is the one narrow format a CPU can execute natively — bf16 has no C type and no general ARM/x86 arithmetic, and stays emulated by design — and where the hardware has it, the lane count doubles against f32.

Off by default, and the asymmetry with narrow_compute_f32 is deliberate: computing in fp16 keeps intermediates at fp16's 10-bit mantissa and 65504 range, so it trades accuracy for throughput, while widening to f32 trades nothing. It also only pays where the target's arithmetic is genuinely 16-bit (Ir.Backend_intf.hardware_limits.native_fp16_arithmetic); on a target that merely promotes to float, turning it on costs accuracy for no speed, so the backend ignores it there.