FP8 Lands in XiangShan's Reference Model: Inside the Verification Stack for RISC-V Matrix Extensions

Every RISC-V matrix extension argument eventually collapses into the same question, and it is never the peak-throughput one. It is this: when two implementations disagree by one unit in the last place, which one is wrong? A matrix unit is a reduction tree with an accumulator, a rounding mode, a saturation flag and a signedness decision per operand pair, and any of those can differ between a vendor's RTL and a software model without either side being obviously broken. Peak TOPS is a marketing number. Agreement on the edge cases is the thing that lets a design tape out.

That is why the more interesting RISC-V news of the past week is not a chip launch. It is seventeen commits landed across OpenXiangShan/NEMU and OpenXiangShan/difftest between 24 and 27 October 2026, which together move the XiangShan project's matrix-extension support from "the instruction semantics are implemented" to "the numerical semantics are verifiable". FP8 arrives in the reference model with FP32 accumulation. MMA event types stop being inferred from storage width and start being decoded from the mcfg register. INT32 saturation becomes real on both the software and the GPU differencing backends. And synchronous verification becomes a build flag.

What the Matrix Extension Actually Looks Like in Code

NEMU gates the whole feature behind CONFIG_RV_AME. The register file is modelled as matrix tile registers plus a set of configuration registers, and the instruction set is small enough to enumerate. Reading src/isa/riscv64/instr/ame/ on 29 October 2026 gives this:

The geometry is compile-time configured. Defaults in src/isa/riscv64/Kconfig are RV_AME_TLEN = 65536 bits (8 KiB per tile register) and RV_AME_TRLEN = 512 bits (64 bytes per row), which yields ROWNUM = TLEN / TRLEN = 128 rows. RV_AME_MLEN = 32 bits drives the accumulator sizing: ALEN = ROWNUM × ROWNUM × MELEN = 64 KiB, and ARLEN = ROWNUM × MELEN = 512 bytes. Tile bounds follow directly: TMMAX = TNMAX = ROWNUM, while TKMAX(sew) = TRLEN / sew — so with a 512-bit row the reduction depth a single mmacc may request is 64 elements at 8-bit, 32 at 16-bit and 16 at 32-bit. Widening the data type costs reduction depth, which is why tile shape and type have to be decided together.

The type encoding is a 15-entry table in mreg.h. It is worth reproducing because three of its entries are the reason a naive verification harness silently passes:

enum MTYPECODE_TAB0 {
  MTYPECODE_INT4 = 0,  MTYPECODE_UINT4,   MTYPECODE_INT8,
  MTYPECODE_UINT8,     MTYPECODE_INT32,   MTYPECODE_NVFP4,
  MTYPECODE_MXFP4,     MTYPECODE_FP8E5M2, MTYPECODE_FP8E4M3,
  MTYPECODE_FP16,      MTYPECODE_BF16,    MTYPECODE_TF32,
  MTYPECODE_FP32,      MTYPECODE_FP2PACK4, MTYPECODE_FP2PACK5
};

INT8 and UINT8 are one byte wide with different arithmetic meaning. FP8E5M2 and FP8E4M3 are one byte wide with different exponent and mantissa splits (5/2 versus 4/3), and INT8 is also one byte wide. Three distinct semantics, one storage width. Any code path that carries only "how many bytes" is already wrong, which is exactly what commit #1269 was written to fix.

The Type Matrix, Before and After #1266

mcompute_impl.h enumerates the multiply-accumulate combinations the model accepts. Comparing 24 October against 27 October gives a clean before/after:

The E4M3/E5M2 split is not cosmetic. E4M3 gives four exponent bits and three mantissa bits — more precision, narrower dynamic range, no infinity encoding. E5M2 gives five and two — range at the cost of mantissa. Training pipelines reach for one and inference stacks for the other, and a reference model that implements "FP8" as a single type will disagree with hardware on exactly the overflow-heavy inputs that matter.

FP4 is conspicuously not in the accepted list, and the reason is in mmacc itself:

#ifdef CONFIG_RV_AME_FP4
  if (is_raw_fp4_type_code(s1mcfg) || is_raw_fp4_type_code(s2mcfg) ||
      is_raw_fp4_type_code(dmcfg)) {
    longjmp_exception(EX_II);
  }
#endif

NVFP4 and MXFP4 exist in the type table but mmacc raises an illegal-instruction exception when it sees them. The encoding knows about FP4; the execute path does not yet implement it. That distinction is the difference between a roadmap and a claim.

Seventeen Commits in Ninety-Six Hours

Querying both repositories for the window 2026-10-24T00:00Z to 2026-10-28T00:00Z returns 7 commits in NEMU and 10 in difftest. Reproduced verbatim, newest first:

Note the shape of it. The 25th is throughput work — Tensor Cores, per-format kernel specialisation, fewer copies. The 26th is correctness work — semantics, flags, trace replay. The 27th is completeness and determinism — a new data type and a switch that trades parallelism for reproducibility. That is a team working through a verification stack in the order you would want them to.

Deep Dive 1: FP8 by Unifying, Not Bolting On

The interesting part of #1266 is the diffstat. mcompute_impl.c goes +104 / −153 — the file shrinks by 49 lines while gaining a data type. The commit message states the mechanism plainly: "Unify floating-point MMA vectorization and extend it to FP8."

The per-file breakdown:

The type-matching logic that results is deliberately boring, which is the point — one branch per accepted combination, and everything else falls through to FLOAT_MMACC_UNSUPPORTED:

if (d_type == MTYPECODE_FP16 &&
    s1_type == MTYPECODE_FP16 && s2_type == MTYPECODE_FP16) {
  return FLOAT_MMACC_FP16_FP16_FP16;
} else if (d_type == MTYPECODE_FP32 &&
    s1_type == MTYPECODE_FP16 && s2_type == MTYPECODE_FP16) {
  return FLOAT_MMACC_FP16_FP16_FP32;
} else if (d_type == MTYPECODE_FP32 &&
    s1_type == MTYPECODE_BF16 && s2_type == MTYPECODE_BF16) {
  return FLOAT_MMACC_BF16_BF16_FP32;
/* ... FP32, then the two new entries ... */
} else if (d_type == MTYPECODE_FP32 &&
    s1_type == MTYPECODE_FP8E4M3 && s2_type == MTYPECODE_FP8E4M3) {
  return FLOAT_MMACC_FP8E4M3_FP8E4M3_FP32;
} else if (d_type == MTYPECODE_FP32 &&
    s1_type == MTYPECODE_FP8E5M2 && s2_type == MTYPECODE_FP8E5M2) {
  return FLOAT_MMACC_FP8E5M2_FP8E5M2_FP32;
}

Mixed FP8 inputs — E4M3 on one side, E5M2 on the other — are not in the accepted set. Neither is FP8 into an FP16 accumulator. Whether hardware should support either is a spec question, not an implementation one; right now the model's answer is no, and a differential test that feeds such a pair should expect an exception rather than a result.

Deep Dive 2: Why mcfg Semantics Beat Storage Width

#1269 is the commit that will save somebody a very bad week. Its message: "Separate AMU type payloads from byte-storage widths and integer arithmetic signedness. Use the semantic encoding for both MMA sources and the destination in REF, CTRL and diagnostic events." Three files, +41 / −7.

The failure mode it closes is subtle. Differencing does not compare operands — it compares events: a compact record saying "an MMA happened at this PC, with these operands, this type, this shape". If that record encodes the operand type as a storage width, then E5M2, E4M3 and INT8 all become "1 byte", and a model that computes E4M3 against RTL that computed INT8 will not be caught — the event descriptor matched. You get a green test and a wrong chip.

The fix moves the encoding to the mcfg register's semantics rather than its size, which means the event now carries the actual type code out of the 15-entry table. It is a one-line concept and an unavoidable one: the differencing ABI must be as typed as the ISA.

Deep Dive 3: Saturation, the Case the Docstring Admitted

Before #1008, the difftest documentation said this of the INT8 CUDA backend:

Integer MMA uses INT8 Tensor Core instructions on SM80 and newer targets,
retaining the software kernel on older targets. Accumulation wraps at 32
bits without saturation.

After #1008 the same paragraph reads:

Integer MMA uses INT8 Tensor Core instructions on SM80 and newer targets,
retaining the software kernel on older targets. Both paths support wrapping
and saturating INT32 accumulation.

One sentence changed; the consequences run through eight files. In the CUDA kernel the previous comment above the PTX was explicit about the choice:

// Before:
// No .satfinite: the accumulator wraps modulo 2^32, as in the integer reference model.

// After:
// Half dots start from zero; the caller applies the selected INT32 overflow mode.

The instruction itself is mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 — an M16 × N8 × K32 INT8 tile with an S32 accumulator, which is the shape most 8-bit inference kernels are built around. What changed is that the sat field now travels with the request instead of being assumed:

The new helper is about as small as code gets and about as easy to get wrong:

CUTE_MMA_HOST_DEVICE static inline int64_t add_int32(int64_t lhs, int64_t rhs, bool saturate) {
  const int64_t result = lhs + rhs;
  if (saturate) {
    if (result > INT32_MAX) return INT32_MAX;
    if (result < INT32_MIN) return INT32_MIN;
    return result;
  }
  return wrap_signed(result, 32);
}

The comment above the calling loop is the tell that this was a real correctness fix and not a feature: "Follow the same ReducePE chunk boundaries for wrapping and saturation." A reduction tree that clamps at the end and one that clamps per chunk produce different answers, and the model has to clamp where the hardware clamps. That is the level at which matrix extensions are actually specified — or, today, not yet specified, which is why the model has to choose and the choice has to be recorded.

Deep Dive 4: Determinism as a Build Flag

Asynchronous differencing is how you get throughput: submit a buffer, let a worker thread verify it while the DUT keeps running, compare later. It is also how you lose error attribution, because by the time a mismatch surfaces the core has moved on and the reported PC may belong to a different instruction.

#1009 adds the escape hatch in cuda.mk, five lines including the comment:

MMA_ASYNC ?= 1
ifeq ($(MMA_ASYNC),0)
SIM_CXXFLAGS += -DCONFIG_DIFFTEST_MMA_SYNC
endif

With MMA_ASYNC=0 the verifier stops owning a thread. The header documents the contract change directly: "With CONFIG_DIFFTEST_MMA_SYNC, submission verifies synchronously on the caller; start()/stop() are no-ops and flush() reports the retained error." The checker change in matrix.cpp is the part that matters for debugging:

+#ifdef CONFIG_DIFFTEST_MMA_SYNC
+          // Do not consume the next ready event after a synchronous failure.
+          if (mma_verifier->has_mma_verification_error()) {
+            set_error_pc(state->coreid, amu_event.pc);
+            return STATE_ERROR;
+          }
+#endif

That is the whole value of the flag: on failure, stop, and attribute the error to the PC that caused it. The default stays asynchronous (MMA_ASYNC ?= 1), with MMA_BATCH_SIZE ?= 256 and MMA_BACKPRESSURE_FACTOR ?= 4. So the same codebase gives you a fast mode for regression runs and a deterministic mode for the afternoon you are chasing a one-in-ten-thousand mismatch.

The Supporting Cast

Three commits outside the MMA path are worth knowing because they change what the harness can do:

• #998, BOOM DiffTest support (26 Oct) — the out-of-order BOOM core gets a differencing path of its own, paired with NEMU's #1262 layered BOOM reference profiles. Matrix verification on an in-order model is a different problem from matrix verification behind a reorder buffer.
• #981, reftrace (26 Oct) — "record and replay REF-driving traces". When a mismatch takes hours to reproduce, being able to replay the exact reference-model input sequence is the difference between a bug report and a bug fix.
• #1006, FPGA split (27 Oct) — separates the DiffTest link from reference-model checking for FPGA targets, so a link-level stall no longer looks like a numerical mismatch.

And two smaller ones that are quietly important: #1268 preserves host FMA exception flags across model execution — without it, floating-point status bits leak between the model and its host and a differencing run can fail for reasons that have nothing to do with the DUT; #1270 reuses the HostTLB for shared FAST references, which is the kind of change that makes a 64-core run finish before lunch.

How to Read This

Three observations, offered as reading rather than conclusion.

First, the bottleneck has moved. A matrix unit with 128-row tiles and a 64 KiB accumulator is not something you can step through cycle by cycle in software any more. The 25 October commits moved INT8 comparison onto Tensor Cores precisely because the software kernel could not keep up with the DUT. Verification infrastructure is becoming hardware-accelerated in its own right, and that is a permanent change in how these projects are staffed.

Second, the unspecifiable parts are being specified by code. Saturation on overflow, the chunking of a reduction tree, whether mixed FP8 inputs are legal, what happens to NaN in E4M3 (which has no infinity encoding at all) — none of these are settled in a ratified document for matrix extensions today. They are being settled by whichever implementation ships first and by whether the reference model matches it. That is a normal phase for an ISA extension and a genuinely risky one for anyone building a second implementation against a document that does not yet say.

Third, the honest gaps are visible in the source. FP4 is encoded and rejected. Mixed FP8 pairs are rejected. There is no ratified spec to cite. A reader who wants to know what RISC-V matrix extensions actually do in October 2026 is better served by mcompute_impl.h than by any slide deck, because the code has to answer the questions the slides defer.

For the ecosystem, that is the useful signal in a week with no chip launch: the open-source RISC-V matrix stack is being hardened at the numerical-semantics layer, in public, with commit messages that explain the trade-off. When the specification does land, the projects that have been differencing against a typed reference model will be the ones that can claim compliance without an asterisk.