Relu fusion

One instruction, three possible semantics

The fused relu is one line in the epilogue:

acc[m] = __riscv_vfmax_vf_f32m1(acc[m], 0.0f, vl);

vfmax performs maximumNumber: when one operand is NaN the result is the other operand, so relu of NaN is 0. That matches C’s fmaxf(x, 0.0f) and the usual x > 0 ? x : 0 form.

MLIR has three ops that look interchangeable and are not. arith.maximumf treats -0.0 as less than +0.0 and propagates NaN. arith.maxnumf returns the non-NaN operand and leaves the signed-zero case open. arith.maximumnumf is the IEEE maximumNumber form, which is what the instruction implements. Pick the one that matches the reference rather than the one that is shortest to type.

arith.maximumf cannot become a single vfmax, because NaN propagation needs a comparison and a select around the max. That is a useful signal: if the epilogue disassembles to one vfmax.vf, you wrote a quieting maximum, and if it is three instructions, you wrote a propagating one.

Cost removed

One more read and one more write of the output per element, 8 bytes for f32, and one more traversal of the tile set. On a fused epilogue the marginal cost of the relu is the instruction that was already in a register.

Order, twice

Bias then relu, because relu(x + b) with a negative b clips a different set of values than relu(x) + b. And relu after the full K loop, for the reason in p2-09: a partial sum is not the value the activation is defined on.

Things that break

arith.maximumf in the IR against an fmaxf reference. Random test data contains no NaN, so the suite passes; the difference appears on real inputs as a NaN written where the reference has zero. Add one test row with a NaN and one with -0.0. The second failure is relu inside the K loop, which agrees at K=1 and clips partial sums from K=2 up, producing output that is low but never obviously wrong.