diff --git a/AGENTS.md b/AGENTS.md index e74c942e25..1324ebf732 100644 --- a/AGENTS.md +++ b/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= --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. diff --git a/VX_config.toml b/VX_config.toml index c7206c6ed5..f3a1d6a238 100644 --- a/VX_config.toml +++ b/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))" diff --git a/ci/baselines/perf/core.json b/ci/baselines/perf/core.json index 263dbd77b7..04d7bddb65 100644 --- a/ci/baselines/perf/core.json +++ b/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", diff --git a/ci/baselines/perf/dxa.json b/ci/baselines/perf/dxa.json index 02d6bb0426..a11cc9611b 100644 --- a/ci/baselines/perf/dxa.json +++ b/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", diff --git a/ci/baselines/perf/graphics.json b/ci/baselines/perf/graphics.json index 80afa627b1..6e35026942 100644 --- a/ci/baselines/perf/graphics.json +++ b/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", diff --git a/ci/baselines/perf/tensor.json b/ci/baselines/perf/tensor.json index 6d75562e72..fb9764766a 100644 --- a/ci/baselines/perf/tensor.json +++ b/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", diff --git a/ci/baselines/perf/tensor_wg.json b/ci/baselines/perf/tensor_wg.json index 09144589da..2e77cfa966 100644 --- a/ci/baselines/perf/tensor_wg.json +++ b/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", diff --git a/ci/testcases/core.yaml b/ci/testcases/core.yaml index 0f736012aa..92532beef2 100644 --- a/ci/testcases/core.yaml +++ b/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 diff --git a/ci/testcases/dxa.yaml b/ci/testcases/dxa.yaml index 8eb10b7441..78e0ee6cba 100644 --- a/ci/testcases/dxa.yaml +++ b/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 diff --git a/docs/debugging.md b/docs/debugging.md index fb8253cbb6..63e0f983d7 100644 --- a/docs/debugging.md +++ b/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 diff --git a/docs/designs/lsu_pipeline_design.md b/docs/designs/lsu_pipeline_design.md index b9a0c3a469..46aafb5b6e 100644 --- a/docs/designs/lsu_pipeline_design.md +++ b/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 diff --git a/hw/rtl/VX_gpu_pkg.sv b/hw/rtl/VX_gpu_pkg.sv index 978e902505..2fe74f0350 100644 --- a/hw/rtl/VX_gpu_pkg.sv +++ b/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 diff --git a/hw/rtl/core/VX_lsu_agu.sv b/hw/rtl/core/VX_lsu_agu.sv index bedd336e39..658903cc92 100644 --- a/hw/rtl/core/VX_lsu_agu.sv +++ b/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 diff --git a/hw/rtl/tcu/VX_tcu_agu.sv b/hw/rtl/tcu/VX_tcu_agu.sv index 087b825c37..e6e5cc6dce 100644 --- a/hw/rtl/tcu/VX_tcu_agu.sv +++ b/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] diff --git a/sim/avedsim/Makefile b/sim/avedsim/Makefile index c237b3db2d..539c3922ef 100644 --- a/sim/avedsim/Makefile +++ b/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 diff --git a/sim/opaesim/Makefile b/sim/opaesim/Makefile index a5ee3298f1..5fc6e6aaee 100644 --- a/sim/opaesim/Makefile +++ b/sim/opaesim/Makefile @@ -8,6 +8,10 @@ AFU_DIR := $(RTL_DIR)/afu/opae 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$(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 diff --git a/sim/rtlsim/Makefile b/sim/rtlsim/Makefile index 8d2ea58b1b..b935e8458a 100644 --- a/sim/rtlsim/Makefile +++ b/sim/rtlsim/Makefile @@ -7,6 +7,10 @@ SRC_DIR = $(VORTEX_HOME)/sim/rtlsim 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$(ROOT_DIR)/sw -I$(ROOT_DIR)/hw -I$(SIM_COMMON_DIR) -I$(SW_COMMON_DIR) CXXFLAGS += -I$(THIRD_PARTY_DIR)/softfloat/source/include CXXFLAGS += -I$(THIRD_PARTY_DIR)/ramulator/ext/spdlog/include diff --git a/sim/simx/dtm/debug_module.cpp b/sim/simx/dtm/debug_module.cpp index 110389ceab..be90ffb4c7 100644 --- a/sim/simx/dtm/debug_module.cpp +++ b/sim/simx/dtm/debug_module.cpp @@ -6,6 +6,7 @@ #include #include #include +#include "stack_interleave.h" namespace { @@ -913,7 +914,9 @@ vortex::Word DebugModule::read_program_memory(vortex::Word addr, size_t size) co } // Read the specified number of bytes uint8_t buffer[8] = {0}; // Max 8 bytes for 64-bit access - ram_->read(buffer, static_cast(addr), size); + for (size_t i = 0; i < size && i < sizeof(buffer); ++i) { + ram_->read(&buffer[i], vortex::StackInterleave::map(uint64_t(addr) + i), 1); + } // Convert to Word based on size vortex::Word value = 0; @@ -934,7 +937,9 @@ void DebugModule::write_program_memory(vortex::Word addr, vortex::Word value, si for (size_t i = 0; i < size && i < sizeof(vortex::Word); ++i) { buffer[i] = static_cast((value >> (i * 8)) & 0xFF); } - ram_->write(buffer, static_cast(addr), size); + for (size_t i = 0; i < size && i < sizeof(buffer); ++i) { + ram_->write(&buffer[i], vortex::StackInterleave::map(uint64_t(addr) + i), 1); + } } vortex::Word DebugModule::direct_read_register(uint16_t regaddr) diff --git a/sim/simx/lsu_unit.cpp b/sim/simx/lsu_unit.cpp index 4a76f1741f..19cb1f0017 100644 --- a/sim/simx/lsu_unit.cpp +++ b/sim/simx/lsu_unit.cpp @@ -22,6 +22,7 @@ #include "local_mem_switch.h" #include "constants.h" #include "mem_block_pool.h" +#include "stack_interleave.h" #include "VX_types.h" #include "VX_config.h" @@ -163,6 +164,15 @@ void LsuUnit::compute_addrs(uint32_t b, instr_trace_t* trace) { // upper bits into the 64-bit address field. e.addr = Word(rs1_data[t].i + (uint64_t)stride * rs2_data[t].u + offset); e.size = data_bytes; + if (StackInterleave::contains(e.addr)) { + // Interleaving is word-granular; a wider access would straddle two threads. + if (e.size > StackInterleave::WORD_SIZE) { + std::cout << "Error: stack access wider than a word: addr=0x" << std::hex + << e.addr << std::dec << ", size=" << e.size << std::endl; + std::abort(); + } + e.addr = StackInterleave::map(e.addr); + } // The datapath has no misalignment support: only a naturally aligned // access is guaranteed to sit inside one memory block, which is what // lets a lane's payload pack into a single block downstream. diff --git a/sim/simx/stack_interleave.h b/sim/simx/stack_interleave.h new file mode 100644 index 0000000000..39b4c6a4e2 --- /dev/null +++ b/sim/simx/stack_interleave.h @@ -0,0 +1,66 @@ +// Copyright © 2019-2023 +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#pragma once + +#include "types.h" +#include "constants.h" + +namespace vortex { + +// 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. +class StackInterleave { +public: +#ifdef VX_CFG_LSU_STACK_INTERLEAVE_ENABLE + static constexpr bool ENABLED = (VX_CFG_NUM_THREADS > 1); +#else + static constexpr bool ENABLED = false; +#endif + + static constexpr uint64_t WORD_SIZE = sizeof(Word); + static constexpr uint64_t STACK_SIZE = uint64_t(1) << VX_MEM_STACK_LOG2_SIZE; + static constexpr uint64_t NUM_THREADS = VX_CFG_NUM_THREADS; + static constexpr uint64_t NUM_STACKS = uint64_t(VX_CFG_NUM_CLUSTERS) * NUM_SOCKETS * VX_CFG_SOCKET_SIZE + * VX_CFG_NUM_WARPS * NUM_THREADS; + static constexpr uint64_t TOP = VX_MEM_STACK_BASE_ADDR; + static constexpr uint64_t SPAN = NUM_STACKS * STACK_SIZE; + static constexpr uint64_t BOTTOM = TOP - SPAN; + + static_assert(ispow2(NUM_THREADS), "stack interleave needs a power-of-two thread count"); + static_assert(SPAN <= TOP, "stack window underflows the address space"); + + static bool contains(uint64_t addr) { + return ENABLED && addr >= BOTTOM && addr < TOP; + } + + static uint64_t map(uint64_t addr) { + if (!contains(addr)) { + return addr; + } + uint64_t off = addr - BOTTOM; + uint64_t stack = off / STACK_SIZE; + uint64_t thread = stack % NUM_THREADS; + uint64_t group = stack / NUM_THREADS; + uint64_t words = STACK_SIZE / WORD_SIZE; + uint64_t word = ((off % STACK_SIZE) / WORD_SIZE) ^ (group % words); + return BOTTOM + ((group * words + word) * NUM_THREADS + thread) * WORD_SIZE + (off % WORD_SIZE); + } +}; + +} diff --git a/sim/simx/tcu/tcu_unit.cpp b/sim/simx/tcu/tcu_unit.cpp index a4513950b1..398aa6ed42 100644 --- a/sim/simx/tcu/tcu_unit.cpp +++ b/sim/simx/tcu/tcu_unit.cpp @@ -21,6 +21,7 @@ #include "local_mem.h" #include "processor_impl.h" #include "mem/memory.h" +#include "stack_interleave.h" #include #include #include @@ -416,6 +417,13 @@ class TcuUnit::Impl { agu.selector = selector; for (uint32_t t = 0; t < VX_CFG_NUM_THREADS; ++t) { agu.addrs.at(t) = base_addr + uint64_t(t) * 4; + // Metadata is read as a linear warp-wide tile, which an interleaved + // per-thread stack cannot provide. + if (StackInterleave::contains(agu.addrs.at(t))) { + std::cout << "Error: TCU metadata inside the thread stack window: addr=0x" + << std::hex << agu.addrs.at(t) << std::dec << std::endl; + std::abort(); + } } } diff --git a/sim/xrtsim/Makefile b/sim/xrtsim/Makefile index d0df3acde7..b2e9be960a 100644 --- a/sim/xrtsim/Makefile +++ b/sim/xrtsim/Makefile @@ -9,6 +9,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$(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