Vector tail handling

Two tails, not one

The vector tail is the row not dividing VLMAX. __riscv_vsetvl_e32m1(N - j) returns min(N - j, VLMAX), so the last operation is short and __riscv_vse32_v_f32m1 writes exactly vl elements. No cleanup loop is needed, and a loop written as for (j = 0; j < N; j += VLMAX) is wrong because it assumes every strip is full.

The tile tail is different: the tile does not divide M or N, so the tile’s own extents run past the shape. That one needs clamped extents, as in affine.min on the tile size, or a predicated access.

Predication

The masked forms take a mask as their first argument, __riscv_vle32_v_f32m1_m(vbool32_t vm, ptr, vl), and at the instruction level the mask lives in v0. Use them when the run you want is not a contiguous prefix of the strip: a strided access, a segment along the wrong dimension, or a tile whose valid region starts partway in.

A shorter vl is cheaper when the valid run is a prefix, which is the common case for a unit-stride N dimension. Masking costs the mask construction, the extra operand, and a v0 that is no longer available to anything else.

Cost shape

A vsetvl tail costs one short vector operation per row. A scalar cleanup loop costs up to VLMAX scalar iterations per row, so its total scales with the number of rows, and for a small M it is most of the kernel. In a sweep the signature is a step down immediately below a VLMAX multiple that recovers as the shape grows; a slope means something else is wrong.

Things that break

vl is computed once for the first strip and reused. The first strip is full, so the second strip starts at the wrong place or runs one element short, and the output is wrong only when N > VLMAX. Compute vl inside the loop from the remaining count. The second failure is an off-by-one cleanup loop with j <= N, which writes one element past the row; in row-major C that lands on the next row’s first element, so the error appears as a wrong value far from the tail.