Fixes and clarifications from the interal review. - #55
Merged
Conversation
… over a guaranteed non-empty set
Two coupled changes to the lambda WARL selection model:
1. Deterministic auto-selection (W1). Where a configuration instruction
must replace the current lambda (a vsetvli/vsetivli, or a zero/
preserve-or-initialize write whose current value is unsupported), it
now selects "the largest supported nonzero lambda value" rather than
an unspecified supported value. Software can therefore predict the
geometry after a SEW change without reading lambda back. The nonzero-
request rule (largest supported <= request, else smallest) was already
deterministic and is unchanged.
2. Guaranteed non-empty supported set (W3 correction). For any legal
(VLEN, SEW) the set of supported nonzero lambda values is non-empty,
because EMUL_C = VLEN / (SEW * lambda^2) <= 16 always admits at least
one power-of-two lambda. The minimum supported lambda is 1 when
VLEN <= 16 * SEW and a larger power of two otherwise. Consequently
lambda[2:0] = 000 is never produced as a *selected* value by a
configuration instruction; it survives only as a write-request
sentinel, as the tile load/store instruction-immediate ("use
vtype.lambda"), and as the as-yet-unconfigured reset state.
Removed accordingly: the "no supported lambda" selected-value case, the
"field is set to 000 when none exists" clauses, and the
"__riscv_vsetlambda established value is zero" / "__riscv_ime_lambda
returns zero when none exists" branches. The defensive
"IME instruction traps when lambda=000" guard is retained but reframed
as covering only the unconfigured vtype.
This reverses the 2026-05-12 decision to drop the lambda=1 guarantee:
the guarantee is now "the supported set is non-empty" (not literally
lambda=1, which is invalid for VLEN > 16 * SEW).
…ambda encoding (clarity) Replace the single-purpose note about the tile load/store instruction-immediate 000 with one consolidated note that enumerates both remaining meanings of the 000 lambda encoding and states it is never a selected value: - as a vtype write request (vsetvl rs2 bits, __riscv_vsetlambda(0)): a preserve-or-initialize request that never writes a literal 000; - as a tile load/store instruction immediate: "use vtype.lambda". This closes the recurring three-way ambiguity (selected-value vs write-request vs instruction-immediate) now that the selected-value "none" case has been removed.
… asymmetry (clarity) Add a note explaining why lambda and bs/altfmt_A/altfmt_B behave differently across a vsetvli/vsetivli: - lambda is geometry-constrained by (VLEN, SEW), so it preserve-or- initializes and always ends up holding a supported nonzero value; - bs/altfmt are not geometry-constrained, so they are retained verbatim and their combination legality is deferred to the consuming IME instruction. The note also states the consequence explicitly: an auto-selected lambda combined with retained bs/altfmt may form an unsupported combination that is not caught at configuration time and is instead rejected (illegal-instruction) by the IME instruction that uses it.
…nt and pressure (clarity) Add a register-placement note to the m16 accumulator-types section: - a sixteen-register group must be sixteen-aligned, so an m16 accumulator occupies v0-v15 or v16-v31 (two placements only); - A and B (each at most m8) must live in the opposite half, so an m16-accumulator kernel consumes the whole register file; - because v0 is the mask register, masked operations are not expressible with an m16 accumulator at v0-v15; this affects only masked tile load/store of the m8 halves, since masked IME multiply-accumulate is not defined.
… register-preserving semantics Replace the flat "bit-preserving register-group operations" description of the m16 pair/unpair pseudo-intrinsics with their distinct cost and aliasing models: - vunpairlo / vunpairhi are register-preserving views: the returned m8 is the low/high eight-register subgroup of the m16 source, no data movement, and the result aliases the source (benign in the usual unpair-then-store pattern, where the source is consumed immediately); - vpair is value-copying: it builds a fresh sixteen-register-aligned m16 and may emit register moves when the inputs are not already the aligned halves; the result is independent of the inputs. This puts the cost on the cold compose/init path (vpair) and keeps the accumulate-then-store hot path (vunpair) free. The no-side-effects list is consolidated into a single closing paragraph.
…ning tile shapes The introduction used the undefined symbol sigma and gave the input panels as (sigma x lambda), which silently assumes W = LMUL = 1. Define sigma = VLEN / (SEW x lambda) at first use, and write the input inner dimension as K_eff = lambda x W x LMUL so the shapes are correct for widening and LMUL > 1 configurations.
…nt of the widening factor The worked example set VL_A = M_tile * K_eff. That is correct only for the non-widening (W=1) kernel shown; the surrounding text invites reuse for the widening variants, where M_tile * K_eff = M_tile * lambda * W * LMUL over-counts the A-tile VL by a factor of W (the A tile is loaded at the C accumulator SEW, with W logical elements packed per SEW-wide storage element). Set VL_A = M_tile * lambda * LMUL_AB (independent of W) and add a comment noting a widening variant reuses the same VL_A unchanged.
…te_single_element SAIL primitives read_single_element and write_single_element are the most-used SAIL primitives (every tile element access goes through them), yet they were never declared, leaving the flat-index -> (register, in-register element) mapping implicit. read_block_scales in particular relies on a 16-bit read from v0 regardless of vtype.SEW. Add commented signature/semantics blocks alongside the other base-style helpers, defining the flat-index addressing (element flat_idx mod epr of register vs + flat_idx / epr, epr = VLEN / EEW) and that the access width is EEW irrespective of vtype.SEW with no numeric conversion.
…rate base altfmt from altfmt_A/altfmt_B Two gaps in the vtype field description: - the base altfmt field (Zvfbfa, output-accumulator format, low bits) was never distinguished from the IME altfmt_A / altfmt_B (input formats, high bits), leaving a reader unable to confirm they do not collide; - no reset value was given for the IME fields, yet the "000 only on an unconfigured vtype" invariant depends on one. State that the IME fields sit immediately below vill at the high end of vtype and are distinct from the Zvfbfa altfmt, and add a note that on reset / vill=1 the IME fields read lambda=000, bs=0, altfmt_A=0, altfmt_B=0. Add an explicit anchor to the altfmt output section so the cross-reference resolves.
… (scale operand is a pair product, not a pure power of two) The fp_internal_scale comment described its operand as an E8M0 scale implementable by pure exponent adjustment. The value actually supplied at the call site is blk_scale = scale_A x scale_B, formed by fp_mul in read_block_scales -- a general fmt_C value that may be inexact or overflow to infinity. Reword the comment so it does not imply the rnd=xct scale step is always exact.
…write (clamp) differs from the instruction immediate (trap) The lambda discovery note suggests programming each encoding and reading back the retained value -- which works on the vtype-write path because an unsupported nonzero request is WARL-clamped. The tile load/store instruction-immediate lambda field behaves oppositely: an unsupported nonzero immediate is not clamped and raises an illegal-instruction exception. Add a note cross-referencing the two opposite failure modes so they are not conflated.
…flicts with v0-v15 placement The partial-column m16 half-store uses a column mask read from v0. When the m16 accumulator is placed at v0-v15, v0 is part of the accumulator group and cannot also be the mask, so a masked half-store is not expressible for that placement. State that software needing masked half-stores must place the m16 accumulator at v16-v31, and that full (unmasked) half-stores work at either placement.
…he non-ISO m16 accumulator types The pair/unpair long-form examples only showed i32 and f32, leaving the naming for the non-ISO m16 accumulator types (OFP8 E4M3/E5M2, BFloat16, etc.) undocumented. State that a pair/unpair triple is defined for every type that has an m16 accumulator type, with the long-form name embedding the same type token as the m16 type, and give representative OFP8/BF16 examples plus the pattern for the remaining widths.
… match the SAIL (transposing default is sigma) The shared memory-layout prose stated a single rs2=x0 default of LD = lambda x LMUL for all four tile instructions. The transposing variants (vmttl.v / vmtts.v) actually default to LD = VLEN / (SEW x lambda) = sigma, both in their own per-instruction prose and in the SAIL (the default changed for the transposing variants but the shared statement was never updated). The two values coincide only at the VLEN/SEW = lambda^2 x LMUL corner, so the blanket statement gave the wrong default for every other transposing transfer. Split the shared statement into the order-preserving case (lambda x LMUL) and the transposing case (VLEN / (SEW x lambda)), matching the per-instruction prose and the SAIL.
…main and require lambda support within it
The earlier WARL change claimed "for any legal (VLEN, SEW) the supported
lambda set is non-empty" and stated the lambda=1 condition as VLEN <= 16
x SEW. Two defects:
- the claim is false for VLEN < SEW (e.g. VLEN=32, SEW=64), a legal RVV
vtype for which no lambda yields EMUL_C in {1,2,4,8,16}; and
- nothing actually required an implementation to support any lambda, so
the deterministic "largest supported" auto-select had no foundation.
Define the IME-legal configuration domain as the (VLEN, SEW) for which at
least one architecturally permissible nonzero lambda exists (EMUL_C in
{1,2,4,8,16}, M >= 1), correct the lambda=1 condition to SEW <= VLEN <=
16 x SEW, and add a normative implementation requirement to support at
least one permissible lambda within the domain. Outside the domain,
lambda stays 000 and IME instructions raise illegal-instruction.
Scope the vsetvli auto-select, the vtype-write preserve-or-initialize and
nonzero-write rules, and the __riscv_vsetlambda contract to the IME-legal
domain accordingly.
… the compute SAIL Every multiply-accumulate Exceptions list mandates that vs1/vs2 form aligned LMUL-register groups and vd forms an aligned EMUL_C-register group, but the operational SAIL never verified it: only the tile load/store paths called check_ime_reg_group, so the compute model would silently accept a misaligned or overrunning accumulator. Add a check_gemm_reg_groups helper (vs1/vs2 against LMUL, vd against EMUL_C, reusing check_ime_reg_group for the alignment-and-bound test) and call it at the entry of all four GEMM dispatchers (int_gemm, fp_gemm, fp_scaled_gemm, int_scaled_gemm), covering all eleven mmacc instructions.
…dening mmacc forms The integer encoding map states that vm=0 is reserved only for vmmacc.vv; for the widening forms vm=0 selects the microscaled integer-MX variant (per the "Integer MX encoding map (vm=0)"). The per-instruction text was inconsistent with this: v8wmmacc.vv listed a "vm = 0 (reserved)" exception bullet, and v8wmmacc/vqmmacc/vwmmacc each carried an "if vm == 0 then Illegal" guard in their Operation SAIL that would reject the very encoding the map defines. Remove the v8wmmacc.vv bullet and drop the SAIL guard from all three widening forms. The W=1 plain forms (vmmacc.vv and vfmmacc.vv), which have no microscaled W=1 variant, keep vm=0 reserved.
… group reduction The disclosed-reduction interface was fp_disclosed_reduce_step, a pairwise fold (S = step(S, prod)) applied product by product. RVBNA -- the specification's own flagship psm=1 algorithm -- is a bulk normalisation: it aligns and sums all g_len x W group terms together and applies a single terminal normalisation. A pairwise fold structurally cannot express that, so the disclosure obligation that underpins the psm=1 reproducibility and certification claims could not actually be met through the interface the spec defined. Replace the fold with fp_disclosed_reduce_group, which receives the whole group of exact products at once (as a list) together with the input formats and the term count (g_len, W). The PSM_DISCLOSED branch now gathers the products and calls the bulk hook exactly once per group. The reduction is defined to be order-independent, and the returned value is still left unrounded to the C format so the existing rnd step applies the terminal rounding (e.g. RVBNA's round-to-odd). Update the two disclosure prose references to the new hook name. The recommended signature included fmt_C; it is omitted here because the group sum is contractually returned before rounding to the C format, and the terminal round is governed by the separate rnd value.
…-product groups The three rnd values were presented as peers, but rnd=frm rounds the partial sum S to the C accumulator format with frm and then the round_frm(C + S) accumulation rounds again. Whenever S is inexact -- which happens for any group holding more than one product (G > 1) -- this double rounds, and can disagree with a single correctly-rounded accumulation. That is the very effect rnd=rto and rnd=xct exist to avoid. Add a note stating that the rnd values are not numerically interchangeable, that rnd=frm double-rounds when G > 1, and that implementations grouping more than one product per partial sum should disclose rnd=rto or rnd=xct; rnd=frm is well-behaved only at G = 1, where S is the exact product and the accumulation is the sole rounding.
…he unscaled path The arithmetic-considerations text claimed that because modular addition is associative and commutative the integer result is uniquely defined regardless of accumulation order or grouping. That holds only for the unscaled integer instructions. The microscaled integer-input forms (vfwimmacc/vfqimmacc/vf8wimmacc) sum each block's products exactly but then convert the block sum to floating point, scale it, and accumulate in a floating-point C -- an accumulation that is non-associative, frm-dependent, and inexact for narrow accumulators (FP16/BF16 cannot hold a wide block sum). Qualify the uniqueness/exactness claim to the unscaled path and add that the microscaled integer path follows the floating-point accumulation model of the Zvvfmm section instead.
…urbed SAIL The tile-load Description prose said inactive and tail destination elements "may be overwritten with 1s" under vtype.vma/vta=1, but the operational SAIL for vmtl.v and vmttl.v writes only active body elements: inactive elements perform no access and tail elements are excluded by the loop bound, leaving both undisturbed. The stores were already consistent (inactive/tail simply do not write memory); only the loads disagreed with their own executable model. Reconcile the prose to the SAIL rather than adding agnostic-fill branches: state that tile loads always leave inactive and tail destination elements undisturbed and that the tail-agnostic and mask-agnostic fill policies do not apply to them. This also matches the undisturbed-without-read tile philosophy already stated for C tiles. Fixes the four load prose sites (the two memory-layout sections and the vmtl.v/vmttl.v Descriptions) and adds an explicit anchor on the C tile tail policy section for the cross-reference.
…er and its fflags rule Two inconsistencies around the microscaling forced-NaN path: - The prose forced NaN on "any decoded scale is NaN", while the SAIL keys on fp_is_NaN(scale_A x scale_B), i.e. the combined block scale. These coincide for E8M0 (NaN only at 0xFF, and NaN propagates through the product) but diverge for any scale format admitting zero and infinity, where 0 x infinity yields a NaN product from two non-NaN scales. - fp_scaled_gemm breaks out of the block loop on the NaN block, so the products and accumulations of all later blocks of that element are skipped along with their exception flags -- contradicting the "fflags = OR over the evaluation of all active output elements" guarantee, which never said what a forced-NaN element's evaluation is. Pin the trigger to the combined block scale in both prose and SAIL (noting the E8M0 and 0 x infinity cases), and define the flag behaviour: a forced-NaN element's evaluation terminates at the NaN-producing block; the scale decode/multiply of that block and everything in earlier blocks accrue flags, while that block's products/accumulation and all later blocks are not evaluated and raise none. State this in the fflags section and annotate the SAIL break so the bitwise-OR guarantee is well-defined.
…ved integer-MX rows Three rows of the integer MX encoding map (vfwimmacc SEW=32 and SEW=64, vfqimmacc SEW=64) are reserved for the entire SEW: altfmt_A, altfmt_B, and altfmt are all wildcards, the C Type is em-dash, and both BS cells are _reserved_. Yet their Input A / Input B columns still named Int16 / Int32, which makes table-driven legality ambiguous -- a tool keying on the input column would read a concrete MX input type for a row that is reserved in every case. Set those input columns to em-dash, matching the other fully-reserved rows in the same table (the sub-byte SEW rows). The partially-reserved rows that name Int8/UInt8 are left untouched: there the altfmt fields are concrete and the type name identifies which signedness combination is reserved, so the name is meaningful.
…mat asymmetry The floating-point encoding map reserves a BF16 accumulator from two FP16 inputs (and an FP16 accumulator from two BF16 inputs) but permits either 16-bit accumulator format when the inputs are mixed FP16/BF16. The table stated this without rationale, leaving an apparent output-legality asymmetry unexplained. Add a note: a same-format input row requires a single 16-bit extension, which defines accumulation only in its own format, so the cross-format output is reserved; a mixed-format row already requires both Zvvfp16mm and Zvvbf16mm, so an implementing part supports both accumulator formats and either output is legal. The asymmetry follows from the per-row extension requirements rather than from an independent output-format rule.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.