Add elementwise pipeline (add/sub/mul/relu, reductions, transcendentals)

[npip99]

Motivation

The current ISA is matmul-only: LOAD / MMA / STORE / barriers. There is no elementwise arithmetic at all — no add, no sub, no mul, no activation, no reduction. The MMA does fp32 FMA internally but only inside the accumulator path; it's not exposed as a general op.

To run actual neural-net workloads (anything beyond D = A @ B), we need at least:

• Bias add (linear layer: D = A @ B + bias)
• Activation (relu / gelu / sigmoid → needs max, exp)
• Normalization (layernorm / softmax → needs sum, max, exp, div, sqrt)
• Positional encoding (RoPE → needs sin, cos)

Without these, every non-matmul step round-trips back to the host, killing the streaming pipeline the architecture is built for.

Scope — tiered

Land in three stages so each tier is independently useful:

Tier 1 — basic ewise + reductions (the load-bearing tier)

• EWISE: tile-wide add, sub, mul, neg, max, min, relu
• REDUCE: tile → scalar sum, max (row-wise and full-tile)
• BROADCAST: scalar → tile (for bias add, layernorm scale/shift)
• Operates on TMEM tiles (read fp32, write fp32). Same async/mbarrier completion model as MMA.

This unlocks: bias-add, ReLU, max-pool, dot-product reductions.

Tier 2 — transcendentals

• EWISE_TRANSCENDENTAL: tile-wide exp, sin, cos, rsqrt (1/sqrt for layernorm)
• DIV: tile-wide reciprocal or full divide

This unlocks: softmax, sigmoid, GELU, layernorm, RoPE, Fourier features.

Tier 3 — fused convenience ops (optional)

• EWISE_FMA: D = a*B + C (avoids two passes for bias-add-then-scale)
• SOFTMAX_ROW: fused row-softmax (max → sub → exp → sum → div)
• Only worth doing if profiling shows the unfused versions are bandwidth-bound

Implementation candidates

Datapath sharing: does Tier 1/2 reuse the MAC array's fp32 adders/multipliers (steal cycles from MMA) or get its own EWISE block beside it?

• Sharing: less area, but MMA throughput drops when ewise is in flight.
• Dedicated: more area (~5-10% of compute_array?), full MMA throughput preserved.
• Probably want dedicated — matches the SFU model in NVIDIA designs.

Transcendentals:

1. LUT + linear interp — small ROM, ~2 cycle, ~1 ULP-class error. Good starting point.
2. CORDIC — iterative, no multiplier, slow (~16 iters). Probably skip.
3. Range-reduction + minimax polynomial — what NVIDIA SFU does. ~5 cycle latency, IEEE-near-correct, most area.

Start with (1) for fp16; revisit (3) if/when fp32 transcendental accuracy becomes a requirement.

Open questions

• Tile-wide or scalar-per-cycle? Tile-wide matches MMA's granularity but blows up area for transcendentals. Probably scalar-per-cycle for transcendentals, tile-wide for tier 1.
• IEEE special-case handling (inf / NaN / subnormal) — fail-fast, saturate, or fully compliant?
• Does REDUCE write back to TMEM (as a 1×N tile) or to a new "scalar register" space? TMEM keeps the model simple.
• Cross-tile reductions (full-matrix sum/max for global softmax-norm) — needed for v1 or deferred?

Why this issue exists now

Pre-implementation request. No code yet. Filing so we have a place to discuss the encoding before we commit to one in ISA.md. Tier 1 should be tackled first — Tier 2 builds on its EWISE framework.