Feature/stack interleave

AGENTS.md

@@ -82,7 +82,7 @@ See [docs/testing.md](docs/testing.md) and [docs/debugging.md](docs/debugging.md
   CONFIGS="-DVX_CFG_EXT_TCU_ENABLE" ./ci/blackbox.sh --driver=simx --app=<app> --threads=8
   ```
 - **`make tests` / `make -C tests/regression` build with *default* macros.** Use `CONFIGS` + explicit per-app rebuild for non-default configurations.
-- **`--rebuild=1` forces a driver rebuild** even if the hardware configuration is unchanged. Use it when iterating on the driver itself; `--rebuild=0` suppresses rebuild regardless.
+- **There is no `--rebuild` option.** `blackbox.sh` re-makes the driver on every run, and make rebuilds it on any source or flag change.
 - **Keep smoke/regression runs lean.** Don't enable debug or perf collection unless the run is explicitly for debug or measurement.
 - **RTL coverage path is `xrt`, not `rtlsim`.** When discussing or planning RTL verification, `xrt` is the canonical path — `rtlsim` bypasses the AFU surface. `rtlsim` remains useful for fast iteration on processor RTL; `xrt` is what proves the full integration.
 - **`ci/regression.sh` is the canonical source of tested configurations.** Use it to discover supported parameter combinations before inventing ad hoc ones.

VX_config.toml

@@ -91,6 +91,9 @@ VX_CFG_NUM_SFU_BLOCKS = 1
 [lsu]
 VX_CFG_NUM_LSU_LANES = "expr: $VX_CFG_SIMD_WIDTH"
 VX_CFG_NUM_LSU_BLOCKS = 1
+# Interleave the per-thread stacks by word, so a warp-wide access to one frame
+# slot fills a single line instead of one same-bank line per thread.
+VX_CFG_LSU_STACK_INTERLEAVE_ENABLE = true
 
 VX_CFG_LSU_LINE_SIZE = "expr: min($VX_CFG_NUM_LSU_LANES * ($VX_CFG_XLEN / 8), $VX_CFG_L1_LINE_SIZE)"
 VX_CFG_LSU_QUEUE_IN_SIZE = "expr: max(4, 2 * ($VX_CFG_SIMD_WIDTH // $VX_CFG_NUM_LSU_LANES))"

ci/baselines/perf/core.json

@@ -1,7 +1,7 @@
 {
   "core:raycast-nt16:rtlsim": {
     "32": {
-      "cycles": 734368,
+      "cycles": 567659,
       "instrs": 137486
     },
     "app": "raycast",
@@ -12,7 +12,7 @@
   },
   "core:raycast-nt4:rtlsim": {
     "32": {
-      "cycles": 1545949,
+      "cycles": 1492954,
       "instrs": 259228
     },
     "app": "raycast",
@@ -67,7 +67,7 @@
   },
   "core:sgemmx-nt16:rtlsim": {
     "32": {
-      "cycles": 1599812,
+      "cycles": 1297311,
       "instrs": 266896
     },
     "app": "sgemmx",
@@ -89,7 +89,7 @@
   },
   "core:sgemv-nt16:rtlsim": {
     "32": {
-      "cycles": 684502,
+      "cycles": 633593,
       "instrs": 79120
     },
     "app": "sgemv",
@@ -100,7 +100,7 @@
   },
   "core:sgemv-nt4:rtlsim": {
     "32": {
-      "cycles": 1095272,
+      "cycles": 1060219,
       "instrs": 316432
     },
     "app": "sgemv",
@@ -111,7 +111,7 @@
   },
   "core:softmax-nt16:rtlsim": {
     "32": {
-      "cycles": 3584289,
+      "cycles": 1274529,
       "instrs": 464089
     },
     "app": "softmax",
@@ -133,7 +133,7 @@
   },
   "core:stencil3d-nt16:rtlsim": {
     "32": {
-      "cycles": 2681071,
+      "cycles": 2384949,
       "instrs": 892944
     },
     "app": "stencil3d",

ci/baselines/perf/dxa.json

@@ -12,7 +12,7 @@
   },
   "dxa:perf_gate-copy-nt16:rtlsim": {
     "32": {
-      "cycles": 10781,
+      "cycles": 2098,
       "instrs": 592
     },
     "app": "dxa_copy",
@@ -23,7 +23,7 @@
   },
   "dxa:perf_gate-copy-nt4:rtlsim": {
     "32": {
-      "cycles": 8159,
+      "cycles": 6844,
       "instrs": 1792
     },
     "app": "dxa_copy",
@@ -56,7 +56,7 @@
   },
   "dxa:perf_gate-sgemm2-dxa-nt16:rtlsim": {
     "32": {
-      "cycles": 4435474,
+      "cycles": 3975334,
       "instrs": 1263632
     },
     "app": "sgemm2_dxa",
@@ -67,7 +67,7 @@
   },
   "dxa:perf_gate-sgemm2-dxa-nt4:rtlsim": {
     "32": {
-      "cycles": 15118837,
+      "cycles": 14746531,
       "instrs": 5017616
     },
     "app": "sgemm2_dxa",
@@ -78,7 +78,7 @@
   },
   "dxa:perf_gate-wgmma-dxa-mcast-nt16:rtlsim": {
     "32": {
-      "cycles": 134782,
+      "cycles": 99299,
       "instrs": 32384
     },
     "app": "sgemm_tcu_wg_dxa_mcast",
@@ -89,7 +89,7 @@
   },
   "dxa:perf_gate-wgmma-dxa-mcast-nt4:rtlsim": {
     "32": {
-      "cycles": 427262,
+      "cycles": 408190,
       "instrs": 234560
     },
     "app": "sgemm_tcu_wg_dxa_mcast",
@@ -100,7 +100,7 @@
   },
   "dxa:perf_gate-wgmma-dxa:rtlsim": {
     "32": {
-      "cycles": 63004,
+      "cycles": 47429,
       "instrs": 10080
     },
     "app": "sgemm_tcu_wg_dxa",
@@ -111,7 +111,7 @@
   },
   "dxa:perf_gate-wgmma-sp-dxa:rtlsim": {
     "32": {
-      "cycles": 74866,
+      "cycles": 68517,
       "instrs": 9328
     },
     "app": "sgemm_tcu_wg_sp_dxa",

ci/baselines/perf/graphics.json

@@ -12,7 +12,7 @@
   },
   "graphics:perf_gate-draw3d-nt16:rtlsim": {
     "32": {
-      "cycles": 203037,
+      "cycles": 149499,
       "instrs": 51475
     },
     "app": "gfx_draw3d",
@@ -34,7 +34,7 @@
   },
   "graphics:perf_gate-om-nt16:rtlsim": {
     "32": {
-      "cycles": 539979,
+      "cycles": 192077,
       "instrs": 69648
     },
     "app": "gfx_om",
@@ -45,7 +45,7 @@
   },
   "graphics:perf_gate-om-nt4:rtlsim": {
     "32": {
-      "cycles": 813751,
+      "cycles": 755028,
       "instrs": 278544
     },
     "app": "gfx_om",
@@ -56,7 +56,7 @@
   },
   "graphics:perf_gate-raster-nt16:rtlsim": {
     "32": {
-      "cycles": 75760,
+      "cycles": 14959,
       "instrs": 4708
     },
     "app": "gfx_raster",
@@ -67,7 +67,7 @@
   },
   "graphics:perf_gate-raster-nt4:rtlsim": {
     "32": {
-      "cycles": 63367,
+      "cycles": 55633,
       "instrs": 18640
     },
     "app": "gfx_raster",
@@ -78,7 +78,7 @@
   },
   "graphics:perf_gate-tex-nt16:rtlsim": {
     "32": {
-      "cycles": 77440,
+      "cycles": 34994,
       "instrs": 12304
     },
     "app": "gfx_tex",

ci/baselines/perf/tensor.json

@@ -1,7 +1,7 @@
 {
   "tensor:perf_gate-fp16-isw4:rtlsim": {
     "32": {
-      "cycles": 204506,
+      "cycles": 195794,
       "instrs": 130592
     },
     "app": "sgemm_tcu",
@@ -12,7 +12,7 @@
   },
   "tensor:perf_gate-fp16-mc-nt16:rtlsim": {
     "32": {
-      "cycles": 59168,
+      "cycles": 44583,
       "instrs": 47296
     },
     "app": "sgemm_tcu",
@@ -34,7 +34,7 @@
   },
   "tensor:perf_gate-fp16:rtlsim": {
     "32": {
-      "cycles": 260319,
+      "cycles": 251208,
       "instrs": 130576
     },
     "app": "sgemm_tcu",

ci/baselines/perf/tensor_wg.json

@@ -12,7 +12,7 @@
   },
   "tensor_wg:perf_gate-wgmma-fp16-ss:rtlsim": {
     "32": {
-      "cycles": 758510,
+      "cycles": 841657,
       "instrs": 525536
     },
     "app": "sgemm_tcu_wg",

ci/testcases/core.yaml

@@ -186,3 +186,10 @@ tests:
   check: model_parity
   via: blackbox
   app: raycast
+# Same raycaster at full warp width: spills and call frames exercise the
+# interleaved stack window with 16 threads per access.
+- id: parity-raycast-nt16
+  check: model_parity
+  via: blackbox
+  app: raycast
+  configs: -DVX_CFG_NUM_THREADS=16 -DVX_CFG_NUM_WARPS=16

ci/testcases/dxa.yaml

@@ -495,11 +495,8 @@ tests:
     threads: [4, 16]
 - id: model_parity-wgmma-dxa
   check: model_parity
-  # Instructions match exactly; cycles diverge ~6% (RTL faster). The gap is
-  # insensitive to every issue/commit/execution-window mechanism the SimX
-  # timing model carries, which points at the DXA producer path's own
-  # timing. Tracked with its wgmma siblings as DXA-path follow-up.
-  known_issue: "SimX over-serializes the DXA-fed warp-group pipeline (instrs match; cycles ~6%)"
+  # Instructions match exactly; cycles agree within ~1.2%. The gap was ~6%
+  # (RTL faster) until the per-thread stacks were interleaved across a warp.
   via: blackbox
   app: sgemm_tcu_wg_dxa
   args: -m 128 -n 128 -k 64

docs/debugging.md

@@ -2,11 +2,9 @@
 
 ## Testing changes to the RTL or simulator GPU driver.
 
-The Blackbox utility script will not pick up your changes if the h/w configuration is the same as the last run.
-To force the utility to build the driver, you need pass the --rebuild=1 option when running tests.
-Using --rebuild=0 will prevent the rebuild even if the h/w configuration is different from last run.
+The Blackbox utility script runs `make` on the driver before every test, so RTL, simulator and configuration changes are picked up automatically; there is nothing to force.
 
-    $ ./ci/blackbox.sh --driver=simx --app=demo --rebuild=1
+    $ ./ci/blackbox.sh --driver=simx --app=demo
 
 ## SimX Debugging
 

docs/designs/lsu_pipeline_design.md

@@ -58,6 +58,20 @@ and packs sub-word stores into the line-width word with the appropriate
 shift (`mem_req_data` shifting + `mem_req_byteen` mask). Misalignment is a
 runtime assertion, not a hardware-handled fault.
 
+**Stack interleave** (`VX_CFG_LSU_STACK_INTERLEAVE_ENABLE`, on by default).
+The kernel ABI gives each hardware thread a contiguous stack
+(`sp = VX_MEM_STACK_BASE_ADDR - hartid << VX_MEM_STACK_LOG2_SIZE`), so the
+same frame slot across a warp's threads sits one stack apart: `NUM_THREADS`
+separate lines, all in the same dcache bank. `VX_lsu_agu` remaps every
+address inside the stack window so that, within each group of `NUM_THREADS`
+stacks, offset `{thread, word, byte}` is stored as `{word ^ group, thread, byte}`
+(word = XLEN bytes). A warp-wide spill then fills one contiguous block, the
+same layout NVIDIA local memory and AMD scratch use. A group spans a power
+of two larger than a cache way, so without the XOR skew every warp's copy of
+a frame slot would land in the same cache set. The remap is a function
+of the address alone, so a pointer into another thread's stack still
+resolves, and software, `sp` and frame offsets are unchanged.
+
 ### 2.2 Fence ordering
 
 `fence_lock` (single bit per slice) is set when a fence's last PID packet

hw/rtl/VX_gpu_pkg.sv

@@ -148,6 +148,12 @@ package VX_gpu_pkg;
     localparam SRC_OPD_WIDTH = `UP(SRC_OPD_BITS);
     localparam NUM_SOCKETS = `UP(`VX_CFG_NUM_CORES / `VX_CFG_SOCKET_SIZE);
 
+    // Per-thread stack window: one stack per hardware thread below the stack base.
+    localparam NUM_STACKS = `VX_CFG_NUM_CLUSTERS * NUM_SOCKETS * `VX_CFG_SOCKET_SIZE * `VX_CFG_NUM_WARPS * `VX_CFG_NUM_THREADS;
+    localparam [`VX_CFG_XLEN-1:0] STACK_WINDOW_TOP    = `VX_CFG_XLEN'(`VX_MEM_STACK_BASE_ADDR);
+    localparam [`VX_CFG_XLEN-1:0] STACK_WINDOW_SPAN   = `VX_CFG_XLEN'(NUM_STACKS) << `VX_MEM_STACK_LOG2_SIZE;
+    localparam [`VX_CFG_XLEN-1:0] STACK_WINDOW_BOTTOM = STACK_WINDOW_TOP - STACK_WINDOW_SPAN;
+
 
     // ===== AMO (RVA) sideband =====================================
     // Fields are only meaningful when amo_valid==1; the bank's reservation

hw/rtl/core/VX_lsu_agu.sv

@@ -19,7 +19,8 @@
 //   - pack-load uop:          addr = base + uop_idx * stride
 //     (uop_idx = offset[1:0], stride = rs2; a 2-bit shift-and-add
 //      collapsed with a 3:2 compressor — no multiplier)
-// Both forms are computed in parallel and selected by `pack`.
+// Both forms are computed in parallel and selected by `pack`. Addresses in the
+// per-thread stack window are then interleaved by word across threads.
 
 module VX_lsu_agu import VX_gpu_pkg::*; (
     input  wire [`VX_CFG_XLEN-1:0] base,    // rs1
@@ -50,6 +51,42 @@ module VX_lsu_agu import VX_gpu_pkg::*; (
     // plain: base + sext(offset)
     wire [`VX_CFG_XLEN-1:0] offset_addr = base + `SEXT(`VX_CFG_XLEN, offset);
 
-    assign addr = is_pack ? pack_addr : offset_addr;
+    wire [`VX_CFG_XLEN-1:0] lin_addr = is_pack ? pack_addr : offset_addr;
+
+`ifdef VX_CFG_LSU_STACK_INTERLEAVE_ENABLE
+    // Thread t's stack is the STACK_SIZE bytes below STACK_TOP - t*STACK_SIZE.
+    // Within each group of NUM_THREADS stacks, offset {thread, word, byte} is
+    // stored as {word ^ group, thread, byte}, so one frame slot across a warp's
+    // threads is contiguous. The XOR keeps equal frame slots of different warps
+    // out of the same cache set, since a group spans a power of two. The map
+    // depends on the address alone, so any thread may dereference another
+    // thread's stack pointer.
+    if (`VX_CFG_NUM_THREADS > 1) begin : g_stack_interleave
+        localparam WORD_BITS   = `CLOG2(`VX_CFG_XLEN / 8);
+        localparam THREAD_BITS = `CLOG2(`VX_CFG_NUM_THREADS);
+        localparam STACK_BITS  = `VX_MEM_STACK_LOG2_SIZE;
+        localparam SLOT_BITS   = STACK_BITS - WORD_BITS;
+        localparam GROUP_BITS  = STACK_BITS + THREAD_BITS;
+        `STATIC_ASSERT(`IS_POW2(`VX_CFG_NUM_THREADS), ("invalid parameter: NUM_THREADS=%0d", `VX_CFG_NUM_THREADS))
+        `STATIC_ASSERT(`VX_CFG_FLEN <= `VX_CFG_XLEN, ("invalid parameter: stack accesses must fit a word"))
+        `STATIC_ASSERT(STACK_WINDOW_SPAN <= STACK_WINDOW_TOP, ("invalid parameter: stack window underflows"))
+
+        wire in_stack = (lin_addr >= STACK_WINDOW_BOTTOM) && (lin_addr < STACK_WINDOW_TOP);
+        wire [`VX_CFG_XLEN-1:0] stack_off = lin_addr - STACK_WINDOW_BOTTOM;
+        wire [`VX_CFG_XLEN-GROUP_BITS-1:0] stack_group = stack_off[`VX_CFG_XLEN-1:GROUP_BITS];
+        wire [SLOT_BITS-1:0] stack_slot = stack_off[STACK_BITS-1:WORD_BITS] ^ SLOT_BITS'(stack_group);
+        wire [`VX_CFG_XLEN-1:0] swz_off = {
+            stack_group,
+            stack_slot,
+            stack_off[GROUP_BITS-1:STACK_BITS],
+            stack_off[WORD_BITS-1:0]
+        };
+        assign addr = in_stack ? (STACK_WINDOW_BOTTOM + swz_off) : lin_addr;
+    end else begin : g_stack_linear
+        assign addr = lin_addr;
+    end
+`else
+    assign addr = lin_addr;
+`endif
 
 endmodule

hw/rtl/tcu/VX_tcu_agu.sv

@@ -253,6 +253,15 @@ module VX_tcu_agu import VX_gpu_pkg::*, VX_tcu_pkg::*; #(
     end
     `UNUSED_PARAM (PWD)
     `STATIC_ASSERT (LSU_CLIENT_TAG_WIDTH >= BLOCK_IDX_BITS, ("LSU client tag cannot encode TCU block"))
+`ifdef VX_CFG_LSU_STACK_INTERLEAVE_ENABLE
+    // Metadata is read as a linear warp-wide tile, which an interleaved
+    // per-thread stack cannot provide.
+    if (`VX_CFG_NUM_THREADS > 1) begin : g_meta_stack_check
+        wire [`VX_CFG_XLEN-1:0] meta_end_addr = base_addr + `VX_CFG_XLEN'(NUM_LANES * 4 - 1);
+        `RUNTIME_ASSERT(~req_fire || (meta_end_addr < STACK_WINDOW_BOTTOM) || (base_addr >= STACK_WINDOW_TOP),
+            ("%t: *** %s TCU metadata inside the thread stack window: addr=0x%0h", $time, INSTANCE_ID, base_addr))
+    end
+`endif
     assign client_if.req_valid  = req_valid;
     assign client_if.req_data   = req_w;
     assign client_if.rsp_ready  = busy_r[rsp_block] && issued_r[rsp_block]

sim/avedsim/Makefile

@@ -12,6 +12,10 @@ AFU_COMMON_DIR := $(RTL_DIR)/afu/common
 
 CXXFLAGS += -std=c++17 -Wall -Wextra -Wfatal-errors -Wno-array-bounds
 CXXFLAGS += -fPIC -Wno-maybe-uninitialized
+# GCC 13 turns the generated per-word copy of a wide signal into a memcpy and
+# inlines it with aligned 16-byte stores, faulting on a destination that is only
+# 8-byte aligned in the root class (NUM_THREADS=16 + RTU crashes in reset).
+CXXFLAGS += -fno-tree-loop-distribute-patterns
 CXXFLAGS += -I$(SRC_DIR) -I$(XRTSIM_DIR) -I$(ROOT_DIR)/sw -I$(ROOT_DIR)/hw -I$(SIM_COMMON_DIR) -I$(SW_COMMON_DIR) -I$(DESTDIR)
 CXXFLAGS += -I/$(THIRD_PARTY_DIR)/softfloat/source/include
 CXXFLAGS += -I$(THIRD_PARTY_DIR)/ramulator/ext/spdlog/include