[tinebp]

Summary

This PR adds memory-level parallelism (MLP) to the memory system, in two phases:

• Phase 1: more requests in flight at every level of the memory system.
• Phase 2: cheaper fences, read merging, and closer agreement between SimX and the RTL.

The design is documented in docs/designs/mlp_optimizations.md.

Perf vs master 50cb174 (rtlsim perf_gate, XLEN 32, CI toolchain LLVM v3.0.1):

• Geomean is −3.36% cycles over all 49 cases.
• Retired instruction counts are identical in all 49 cases.

What changes

Phase 1: requests in flight (afd4331 and earlier)

• MSHR line table, so a cache can hold more misses in flight.
• Hashed bank select, so power-of-two strides spread across banks.
• Back-to-back fills into the bank pipeline.
• Write-allocate without fetching the line at a write-back LLC.
• Flush:
  • one broadcast write, with per-port flush tokens;
  • only the L1s flushed between launches when the LLC is the coherence point.
• Fences:
  • stop at the point of coherence;
  • IO fences skip the cache walk.
• Each LSU block has its own memory path.
• Each TCU block has its own metadata port and its own TCU_LD AGU.
• AFU: one AXI port per GPU memory port, with many transactions in flight.
• The command processor's DMA goes through the shared LLC.
• Fetch width is decoupled from issue width (VX_CFG_FETCH_WIDTH).
• Core, socket and cluster counts no longer have to be powers of two.
• Redesigned DRAM model (independent channels, refresh, configured from data-sheet values).
• Hardware profiles: a100, h100, v100, rtx1080, tinygpu.
• RV64 portability fixes across sw/tests.
• Zero-warnings build.
• SimX:
  • matches the RTL's buffer structure, arbiters and registered ports at each pipeline boundary;
  • follows the RTL in KMU, CTA retirement, barrier and collector timing.

Phase 2: fences, merging, parity (347c914..c08c245)

• Memory fences:
  • split into acquire and release;
  • .aq/.rl AMOs are ordered;
  • .aq/.rl fences are dropped at a coherent dcache.
• Dcache read merge:
  • a warp's identical loads are served with one bank access;
  • port 0 now joins the merge.
• Same-line icache fetches from a core's fetch slices are merged.
• Lazy-fetch:
  • a fully written sector is refetched without a memory read;
  • fixed a dcache deadlock between AMO write-back and fills.
• A draining write-through chain stays pending; bank paths retimed.
• Kernel args are staged in clean slots, without a cache flush.
• TCU output is fp32 only, with bit-exact TFR numerics in SimX, plus a tcu_unit test.
• SimX:
  • models the miss path's buffers, LSU ingress, SFU result buffers and the lmem crossbar as the RTL does;
  • keeps the LLC across a launch reset.

Since merging master (058d0e7..40ec5f4)

• Merges origin/master 50cb174, which brings in Rv64 fixes #427, dispatch fix on single-packet inst #428, MUFU + Zfhmin #432 (MUFU and Zfhmin) and RTU performance counters + perf-counter fixes #435 (RTU perf counters).
• Merge fixes:
  • removes a stray MUFU instance from VX_sfu_unit;
  • restores the RTU perf-counter read in VX_csr_mpm (lint failed with PERF_ENABLE and the RTU on).
• SimX RVC fix: an instruction that straddles two fetch words keeps its low half until decode accepts it. Before this, a full decode latch made it decode the wrong half-word and hang. It was seen at RV64 with C on.
• CI: the model_parity and perf_gate cells are split into shards of about an hour each.
• perf_gate baselines re-recorded after the merge. Cycle and instruction counts are unchanged; only the config hashes moved.

Validation

• Final gate on Phase 1 + Phase 2:
  • lint, x32 and x64: 40/40 clean;
  • regression: simx 30, simx xlen64 30, rtlsim 32;
  • cache: simx 27, rtlsim 36;
  • unittest: 10;
  • model_parity: 47 passed. The worst cases are sgemmx at 3.40% and sgemv at 3.38%; the tolerance is 5%.
  • perf_gate: 49 cases recorded.
• After the merge:
  • perf_gate: 49/49 cases have the same cycles and instruction counts as before the merge.
  • perf-enabled lint: 7/7.
  • model_parity for raytracing and MUFU: 4/4.
  • RVC 64-bit regression passes on SimX.
  • perf_gate with RVC on by default: 49/49 cases complete, with no hangs.
• FPGA gate (U55C, 300 MHz target), at c08c245:
  • core: 291.5 MHz;
  • top: 267.9 MHz, against a master baseline built at a 250 MHz target.
• ASIC gate (yosys), at c08c245 plus the catalog fix: all 12 blocks pass; top runs at 618 MHz against a 621 MHz baseline, with +0.5% area.
• ASIC gate at 40ec5f4, on the blocks the merge touched: all pass.
  • core: 586.6 MHz (baseline 586.7), cell area −0.06%.
  • mufu: 1039.5 MHz, identical to the baseline.
  • rtu: 904.5 MHz (baseline 903.4), cell area +0.01%.
• FPGA gate: not rerun since the merge. The RTL the merge changed is mostly off in the default configuration (MUFU, Zfhmin, PERF/RTU counters). The always-on part is a few LUTs, away from the core's critical path.

Known items

• perf_gate cases slower than master: six cases are slower by 0.24–0.57%. The worst is tensor:perf_gate-sgemm2-fp16 at +0.57%. The rest are below 0.4%: copy-mcast, vecadd-nt4, mxfp8, wgmma-fp16-ss and om-nt16.
• sgemm2-fp16: it is sensitive to interleaving. With read merge on and off, the final tree differs by only 180 cycles.
• tensor:perf_gate-fp16-mc-nt16: it is −8.6% vs master but +6.3% vs Phase 1. Port 0 joining the read merge changes how CTAs split across the two cores, and this shape is bound by the slower core.
• Master's own nightly FPGA gate has failed since Oct 17 on the cache build (301.8 → 280.3 MHz). This branch keeps its own cache RTL, which the merge did not change.

🤖 Generated with Claude Code