Feature/stack interleave

[tinebp]

Summary

Vortex gives each hardware thread a contiguous 8 KB stack, so the same stack slot across a warp's threads lands in NUM_THREADS separate cache lines, all in the same dcache bank. Every register spill, callee-saved save/restore and runtime-indexed local becomes NT serialized same-bank misses. With RISC-V's 32 registers, kernels spill a lot, so this cost is significant.

This PR interleaves the per-thread stacks by word in the LSU address generator. Within each group of NUM_THREADS stacks, offset {thread, word, byte} is stored as {word ^ group, thread, byte}. A warp-wide access to one stack slot then fills a single line; NVIDIA local memory and AMD scratch lay out their stacks the same way.

• The XOR skew stops equal slots of different warps from aliasing into one cache set.
• The map is a bijection that depends only on the address, so software, sp, frame offsets and cross-thread stack pointers are unchanged.

Commits

1. lsu: interleave per-thread stacks across a warp's threads
   • VX_CFG_LSU_STACK_INTERLEAVE_ENABLE, default on.
   • RTL is in VX_lsu_agu.sv; SimX is in stack_interleave.h, kept in lockstep with the RTL.
   • TCU metadata inside the stack window is asserted against on both sides.
   • New parity case parity-raycast-nt16.
2. sim: stop GCC 13 miscompiling Verilator models. rt_raycast-nt16 segfaulted during reset, on master too.
   • GCC 13.3 converts the generated per-word copy of a wide signal into memcpy. It then inlines that copy with aligned 16-byte stores (movaps) to an address that is only 8-byte aligned.
   • The fix builds all Verilator models with -fno-tree-loop-distribute-patterns.
3. docs: drop the nonexistent blackbox --rebuild option. AGENTS.md and docs/debugging.md documented it, but blackbox.sh rejects it.

Performance

raycast, 16 warps × 16 threads:

Driver	Before (cycles)	After (cycles)	Speedup
rtlsim	560,672	368,919	1.52x
simx	632,453	381,214	1.66x

24 perf-gate cases improved beyond the 2% ratchet. Their baselines were regenerated with --update-baselines. Examples (rtlsim cycles):

Case	Before	After
softmax-nt16	3,584,289	1,274,529
om-nt16	539,979	192,077
raster-nt16	75,760	14,959
dxa copy-nt16	10,781	2,098
raycast-nt16	734,368	567,659

model_parity-wgmma-dxa's SimX/RTL gap closed from ~6% to 1.2%, so its known_issue is retired.

wgmma-fp16-ss (+11%, baseline regenerated). This is not a cost of the remap: the kernel touches the stack once per warp. Its RTL schedule is bimodal, either ~5.9k cycles per k-step (39% scheduler idle) or ~6.5k (45% idle), and small timing changes flip it.

K sweep at 128×128, rtlsim cycles:

K	master	this PR
32	202,416	199,829
64	389,817	388,047
96	574,918	575,776
128	757,675	841,657
160	944,889	940,841
192	1,253,345	1,130,145
256	1,641,902	1,658,985

Master is in the slow mode at K=192 and 256. This PR moves the slow point to K=128 and out of K=192. The geomean across the sweep is −0.1%. The bimodality exists on master and deserves its own investigation.

Test plan

• pytest ci -m "model_parity or perf_gate" (rtlsim, xlen 32): 79 cases, 76 passed, 1 passed after its known_issue was retired. The other 2 are pre-existing DXA xfails.
• Edge configs (raycast, both drivers, matching instruction counts): NT=1, SOCKET_SIZE=2 with 1 core, 2 cores + L2, RV64 NT=4, RV64 NT=16.
• rt_raycast-nt16 runs again: 52,488 cycles against a baseline of 51,652.