Repository navigation
Rv64 fixes - #427
Merged
Merged
Rv64 fixes#427
Conversation
LinkedList::push_back inserted before the current tail, because insert() places a node ahead of its position: pushing 1, 2, 3 produced 2, 3, 1. The timing wheel and the zero-delay event list are LinkedLists, so two packets sent into one SimChannel in the same cycle with the same delay arrived reversed. Models almost never send twice per channel per cycle, which kept it hidden; a two-wide fetch prototype exposed it as a younger instruction issuing ahead of its producer. push_back now appends at end(). The header also includes <cstddef> so it compiles on its own. tests/unittest/linked_list checks the order and fails against the old header. SimX legs of all 31 model_parity cases: identical retired instructions; cycles move by at most 0.7%, except rt_raycast +1.7%, toward rtlsim (gap 2.4% -> 0.8% against a current rtlsim run). rtlsim and the perf gate do not use this code. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
On RV64 an index computed in uint32_t wraps at 32 bits, so LLVM cannot widen it into a 64-bit induction variable or fold it into a pointer increment, and re-emits zero-extend + shift + add for every access (sgemm's unrolled body: 90 instructions on RV64 vs 35 on RV32). Indices that feed pointer arithmetic are now size_t, which on RV32 is the same type as uint32_t, so RV32 kernels compile to identical code; vm_stress uses int, where size_t moved the loop exit onto a divergent per-thread pointer compare. RV64 correctness: - vx_dxa.h: keep the wgather result (lane 0 carries the shared-memory address) at register width instead of truncating it to 32 bits. - vx_graphics.h: compute the OM aperture address in uintptr_t and pass it to the export at register width; on RV32 the old expression was a 64-bit sum whose dead carry cost an instruction per fragment. - vx_intrinsics.h, vx_tensor.h, packld_box: pack-load strides and ldm steps are size_t; the address unit reads the full register. - vx_serial.S: keep sp 16-byte aligned across the callback (frames 32/64 bytes on RV32/RV64); select the XLEN path with __riscv_xlen, as LOAD_IMMEDIATE64 in src/common.h now does. - vx_start.S: __dso_handle is pointer-sized on RV64. - vx_syscalls.c: use newlib's prototypes for _lseek/_read/_write/_sbrk. - dotproduct2: don't pass uninitialized upper bytes through the shuffle. Kernels whose RV64 cycles got worse with a wider index are left as they were: sgemv, demo, probe, fence, dogfood, dotproduct, softmax, vm_test, mpi_dotproduct, mpi_put_dotproduct, sgemm_tcu, sgemm_tcu_mx, sgemm_tcu_sp_mx and sgemm_tcu_wg. In most of them LLVM's loop strength reduction moves the loop exit onto a per-thread pointer compare, which needs predication every iteration. RV64 simx cycles: sgemm -55%, mstress -28%, sgemm_tcu_wg_mx -22%, sgemm2_dxa -21%, stencil3d -20%, sgemm_v1 -18%, sgemm2_v1 -16%, dxa_copy_mcast -16%, sgemm2 -12%, sgemm2_dxa_mcast -12%, jacobi -10%, sgemm2_tcu -6%; no measured kernel is slower (mpi_sgemm, mpi_blocked_sgemm and the graphics pipeline kernels print no cycle counts and were checked for correctness only). On RV32 only the OM export path and vx_serial's frame change; the affected graphics perf-gate baselines are regenerated with --update-baselines (om -1.3%/-1.9%, draw3d -0.4% cycles on rtlsim). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
On RV64 the trace lost the upper half of both warp-uniform pointers. SimX narrowed the scene base to 32 bits, so a scene above 4 GiB was read from the wrong address and every ray missed. SimX and the RTL both carried the payload pointer as 32 bits and wrote it back through the single VX_RT_PAYLOAD_PTR_LO slot, so a callback shader got a truncated payload pointer. vx_rt_wtrace took uint32_t pointers and narrowed the gathered config, and every tests/raytracing kernel masked its device addresses to 32 bits, which kept all of it hidden: device buffers normally sit below 4 GiB. - The high word goes in slot VX_RT_PAYLOAD_PTR_LO + 1, which already sat unused next to the low word. It is not a new VX_types define: the perf-gate fingerprint hashes VX_types, so a define would mark every perf baseline stale without moving a cycle. - SimX keeps the scene base at the device address width and the payload at full width, and writes the high word on RV64. - RTL widens payload_ptr to XLEN across the RTU bus, the arm staging RAM and the slot descriptor. On RV64 the record write-back sends one more word (the high half) before the hit attribute; STATUS stays last. At XLEN=32 the logic is unchanged. - vx_raytrace.h takes uintptr_t pointers, keeps the gathered config at register width, and adds vx_rt_get_payload(), which reads both words on RV64. - The tests/raytracing kernels pass full-width pointers. On a 64-bit device rt_smoke_chs places its scene and payload above 4 GiB; it fails before this change (miss, payload never written) and passes after. - ci/testcases/raytracing.yaml adds rt_suite-x64, the suite on RV64 simx. - hw/unittest gfx and rtu_core tops follow the wider arm word. RV32: all 34 ray-tracing kernels compile to identical code; the three rt_raycast perf_gate cases reproduce their baseline cycles and instruction counts exactly, the rtlsim suite and model_parity pass, and every perf_gate config hash is unchanged. RV64: the simx suite passes 34/34, and rtlsim passes 31/32 plus rt_smoke_chs above 4 GiB. rt_raycast fails pixel comparison on RV64 rtlsim identically before and after this change (210/768 pixels); that is a separate, pre-existing RTL defect, so rt_suite-x64 stays simx-only until it is fixed. The extra RV64 write-back word moves RV64 rtlsim rt_raycast cycles by 0% to +5.5% depending on shape (NT4 +5.1%, NT16 0.0%, 2-core +0.7%, 8x8 +5.5%) against two added cycles per record; RV64 rtlsim has no perf gate. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
On RV64 with the D extension an FP register is 64 bits, and a single-precision value in it must have its upper word all ones; anything else reads back as the canonical NaN. The core-side RTU unit zero-extended every slot read, so the hit t/u/v that vx_rt_wait (GETWF) lands in f0..f2 arrived as NaN, and the shading that used them went wrong on every hit. rt_raycast failed pixel comparison on RV64 rtlsim (210/768 pixels, maxdiff 172) while simx passed: SimX already NaN-boxes GETWF, as the LSU does for flw and the TCU for its results. The result path now NaN-boxes when the destination is an FP register, the same test the LSU uses; GP reads and RV32 are unchanged. rt_suite-x64 now also runs on rtlsim. RV64 rtlsim: rt_raycast passes with 0/768 mismatched pixels at NT4, NT16, 2-core and 8x8, and its retired instruction count (46136 at NT4) now matches simx. The ray-tracing CI category passes at xlen 64 on simx and rtlsim, and at xlen 32 on rtlsim (suite, model_parity, perf_gate unchanged). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ange Setup converts the screen-space depth plane Z = A'x + B'y + C' to Q7.24 with a plain float-to-int cast. C' is the plane at the screen origin, so a steep plane on a triangle away from the origin puts it past Q7.24's +-128 range even though the depth at every covered pixel is in range. That conversion is undefined: RISC-V fcvt saturates (wrong depth on the device), x86 cvttss2si returns INT_MIN. gfx_pipeline_kernel failed on primbuf[41] at both xlen 32 and 64 with dev C' = INT_MAX against ref C' = INT_MIN. Late-Z and the raster early-Z (SimX and RTL) all evaluate the plane mod 2^32, so a term stored modulo 2^32 gives the exact depth at the covered pixels. PlaneToFixed keeps the plain cast for in-range terms, wraps an out-of-range term through int64, and maps a non-finite or absurdly steep term (past 2^39, which would overflow int64) to 0 so every input is defined on host and device. gfx_setup_kernel, gfx_binsort_kernel and gfx_pipeline_kernel compare the device front end byte for byte against the host reference but were in no CI category; they now run in graphics at xlen 32 and 64. gfx_pipeline_kernel passes at xlen 32 and 64 (failed at both before). graphics simx 38/38 at xlen 32 and 64, vulkan simx 13/13 at xlen 32 and 64, graphics rtlsim 56/56 at xlen 32 including every perf_gate case unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…detour vx_printf and vx_vprintf carried no format attribute, so a kernel could pass a 64-bit value, a pointer or a size_t under a %d and the compiler said nothing; on RV64 that reads the wrong varargs slot. They are now declared format(printf, ...), so -Wall -Werror rejects a mismatched call. tinyprintf already implements the l/ll/z/j/t/h modifiers and %p the checker asks for. Every current caller (vx_spawn.c, tests/kernel/*, tests/regression/printf) is clean at xlen 32 and 64; a deliberately mismatched call is rejected at both. vx_putint copied its int argument into a float before itoa, so any value past 2^24 printed rounded (16777217 printed as 16777216). It now stays an int. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…W_LOAD_B The SW_LOAD_B cooperative load still computed its indices in uint32_t, the pattern the RV64 index-portability change replaced in the SW_LOAD_A loop next to it; it now uses size_t (the same type as uint32_t on RV32). RV64 simx at the dxa sgemm_tcu_wg_dxa-1 config, 8 warps x 32 threads: SW_LOAD_A + SW_LOAD_B 28832 vs 29141 cycles (-1.1%), SW_LOAD_B alone 23167 vs 23160 (flat). With both loads in software no DXA transfer is issued, is_dxa_warp goes unused, and -Werror rejected the build at xlen 32 and 64; it is now [[maybe_unused]], like the descriptor ids above it. RV32 code is identical with SW_LOAD_B, and the default (DXA) build CI runs is identical at xlen 32 and 64. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…m_tcu These five kernels kept uint32_t indices in cbd9cbe because widening them made RV64 slower: LLVM's LoopTermFold moved each loop's uniform exit test onto a per-thread pointer, which needs divergence predication every iteration. llvm_vortex a81a5957f406 keeps a uniform exit uniform on targets with branch divergence, and with it the size_t form is the faster one on RV64 simx: demo -20.4%, fence -20.6%, dogfood -14.6%, softmax -10.7%, sgemm_tcu -8.6% cycles. RV32 code is identical (size_t is uint32_t there). This depends on that compiler change: with the current toolchain the same kernels are 1.4-3.8% slower on RV64 (demo +1.8%, fence +1.4%, dogfood +1.5%, softmax +3.8%, sgemm_tcu +2.1%). It must land after the toolchain prebuilt carrying a81a5957f406, together with the perf baselines that compiler moves. The other kernels left narrow in cbd9cbe stay so: sgemv (+5.0%), dotproduct (+1.5%), sgemm_tcu_mx (+3.7%), sgemm_tcu_sp_mx (+1.8%) and sgemm_tcu_wg (+1.3%) are still slower widened with the new compiler, and vm_test is unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Brings the branch up to master 7fe1936 so it can be proposed as one pull request. Conflicts: - vx_dxa.h: master's new __VX_DXA_ISSUE1/2 macros pinned the gathered operands in uint32_t registers, which re-truncates the DXA shared-memory address on RV64; the macros are kept and widened to size_t. - rt_smoke_ahs_geom: master's capture-buffer rename over the full-width scene/payload pointers. - ci/baselines/perf/graphics.json: master's file; the entries this branch moves are re-recorded with --update-baselines in a later commit. The merged tree equals the rv64-portability-fixes rebase that was validated on master, plus the five widened kernels. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Code merged from master after the RTU pointer widening narrowed the scene again on RV64: - vx_rt_wtrace pinned the gathered trace config in a uint32_t x31 register, truncating the scene base it had just gathered at full width; - the SimX ray log stored scene_root in 32 bits, so a snapshot of a scene above 4 GiB read the wrong memory, and rt_replay replayed it against that address; - rt_smoke_ahs_multi masked its scene pointer to the low word. The ray-log record keeps its 112-byte size (scene_root takes one reserved word); the format version goes from 1 to 2 since the field offsets change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A profile (profiles/<name>.json) names a complete machine: per-core knobs, chip-level knobs, the platform around it, the measurements it is calibrated against, and the sources of every value. gen_config.py layers it between the VX_config.toml defaults and explicit -D flags, validates each knob against the TOML (unknown knobs, enum values, bool/int types, required sources) and folds it into CONFIGS with --merge-configs, one -D per macro so -Werror builds never see a redefinition. Profiles inherit with "extends". config.mk does the merge once, so every flow that includes it -- simulators, runtimes, kernels, tests and the synthesis Makefiles -- builds the same machine without per-flow changes. configure --profile sets a build directory's default and bakes it into the generated headers and the config stamp; blackbox.sh --profile selects one per run. Without a profile the generated flags and headers are byte-identical to before. Adds profiles/rtx_5080.json (GB203 at base clock) with the knobs the model has today. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Warnings had become non-fatal again: a June RTU change re-added Verilator -Wno-fatal to rtlsim after it had been removed, the AFU simulators copied it, and several C/C++ trees (OpenCL/HIP host code, unit tests, kernel library) never had -Werror. Warnings then accumulated unseen. Every C/C++ build (simx, rtlsim, xrtsim, opaesim, avedsim, runtime, kernel library, hw/sw unit tests, OpenCL/HIP/Vulkan apps) is now -Werror and Verilator warnings are fatal, with the code fixed rather than silenced: - RTL: VX_find_first, VX_onehot_encoder (MODEL1) and VX_pipe_buffer rewritten loop-free so Verilator -O0 does not report UNOPTFLAT (each proven equivalent to the original with an sv2v+yosys miter, with negative controls); width fixes in the tex bus arbiter, RTU scheduler config, PLRU replacement, the no-F CSR data path and the TFR MX scale; unused cluster/AFU ports removed or tied off; the early-Z y-pitch add widened explicitly (WIDTHEXPAND), and the RTU scheduler's issue block reads the memory grant instead of mem_fire, which is derived from the block's own mem_issue output (UNOPTFLAT). - Host code: 21 OpenCL apps, HIP, Vulkan, the kernel library and the DXA/FMA testbenches; -Wno-maybe-uninitialized/-Wno-array-bounds dropped; MAKEFLAGS cleared around verilator --build (jobserver warning). - vx_spawn.c printed uint32_t with %d, which the vx_printf format attribute rejects under GCC's newlib (uint32_t is unsigned long there). - The riscv-tests benchmark patch cast the unsuffixed COUT address straight to a pointer (-Wint-to-pointer-cast on RV64), and the riscv-tests stamp did not depend on the patch, so an edited patch was never re-applied. Remaining exceptions compile code that is not ours: the yosys gate-level netlist sims, the cln28hpc vendor macro testbench, and the FireSim driver, whose upstream bridge and platform objects share our CXXFLAGS. Verified: from-clean build of every group at XLEN 32 and 64 with 0 warnings; all 252 CI Verilator commands verilate with 0 warnings; the CI matrix cells run so far (perf_gate, amo, cache, config1 on rtlsim/simx) build warning-free. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A bare -D<ENUM>_<value> override was recognised by splitting the flag at its last underscore, so a value that itself contains one (e.g. -DVX_CFG_PLATFORM_MEMORY_TYPE_HBM2_A100) never matched its enum: the generator silently kept the default, and the profile merge treated the flag and the profile's entry as different knobs and passed both. The enum is now matched by name prefix and the rest checked against its declared values, in both places. Output is byte-identical to before for all 247 CONFIGS strings in ci/testcases at both XLENs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
GPU_DEV_CAPS encoded threads, warps, socket/cluster sizes and clusters as 3-bit log2 fields, so a machine that is not a power of two (48 warps, 84 or 108 cores) reported the wrong shape to the runtime. The spare bits [63:42] now carry exact warps (7 bits), sockets per cluster (8) and clusters (7), produced identically by the RTL CP regfile and the C++ CP model. The decoder prefers them and falls back to the log2 fields only when they read zero, so devices built before this change keep reporting correctly. Verified: every simx/rtlsim test in the matrix so far decodes these caps (NUM_WARPS/NUM_CORES drive the kernel launch), and the A100-shaped 108-core SimX run reports 108 cores. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
With 4 KB pages a working set beyond the TLB reach thrashes: sgemm at N=512 on a VM-enabled RTX 5080-shaped machine ran 18x slower (1% L1 TLB hit rate, the memory port blocked behind 2 requests parked per in-flight page). GPU drivers back allocations with 2 MB pages. VORTEX_VM_MEGAPAGE_HEAP=1 maps the whole heap once with superpages at its identity VA, which the page walker and the L2 TLB's megapage array already support, so a buffer costs no TLB entry per 4 KB (in that study VM then cost 1-4% instead of 18x). Per-buffer access flags are not enforced in this mode and it cannot be combined with VORTEX_RANDOMIZE_VA; the default mapping is unchanged. Adds vm-megapage-heap (sgemm, simx + rtlsim). Verified: vm-megapage-heap passes on simx (XLEN 64); the rtlsim variant runs in the CI matrix vm cell. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Cache::Config stored the input count and the pipeline latency in uint8_t, so any machine whose L2 sees more than 255 inputs (an 84-132 core chip: 16 L1 requests per socket) or whose L2 latency exceeds 255 cycles failed to build with a narrowing error. The RTL caches take both as unbounded parameters; the model now matches. Verified: the 84-core rtx5080 profile builds and runs vecadd on simx; the a100/h100/v100/rtx1080 profiles build (1,728 L2 inputs on a100). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
With local memory disabled (e.g. the tinygpu profile) LMEM_BLK_BASE was never read, and Verilator's UNUSEDPARAM warning, now fatal, stopped the rtlsim build. No CI configuration disables LMEM, so the zero-warnings sweep did not reach it. The window parameters and their static check now sit under the same VX_CFG_LMEM_ENABLE guard as their only use. Verified: 1 core, 2 warps x 2 threads, LMEM disabled builds warning-free on rtlsim and passes sgemm -n16 (15,240 instructions, as on simx). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
vx_device_max_occupancy_grid sizes a 1D block to a whole core and rounds the
grid up, so grid x block exceeds the problem size whenever a core has more
threads than there are points. These four kernels indexed with the global
id unguarded, so the surplus threads read and wrote past their buffers. On a
core with 1,024 threads vecadd -n256 overwrote its kernel-argument block and
later warps faulted on a float loaded as the destination pointer
("misaligned memory access", every app on the 20-132 core profiles). Each
kernel now returns when the index is past num_points, the size the host
already passes, as any kernel on a rounded-up grid must.
The other tests that launch through the helper were audited: they guard the
index, use a grid-stride loop, or size the problem as an exact multiple of
the block.
Verified: the four pass on simx at the default config and at 32 warps x 32
threads, including sizes that do not divide the block (vecadd -n64/-n100/
-n256/-n16384, relu -n64/-n100, dropout -n1000/-n1024, minimal -n1); the
unguarded vecadd reproduces the fault at 32x32 -n256. The guard adds
instructions, so the perf-gate baselines are re-recorded after this series.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The LLC AMO commit engine queues each result in a 4-entry writeback queue that drains through the bank's core-request slot, last in the input priority. Its sizing assumed commit_busy serializes AMOs, but commit_busy only gates core requests: a burst of fills for different lines, each line's chain ending in one AMO, replays an AMO per fill while the fill/replay stream holds the input slot, so nothing drains and the fifth result overflows the queue (amo:zacas:rtlsim, exposed by the per-channel DRAM model's tighter fill bursts; latent on master). - The engine counts the AMOs still able to reach the queue (S0, S0..stC, the compute stage) and flags when the queue cannot take one more; the bank admits a fill only while it can. A fill is admitted only between replay chains and a chain carries at most one AMO, so replays never wait mid-chain, and the drain still fires between chains as before. The flag is registered (the sum grows by at most one a cycle, so it tests one entry early). - req_input_defer missed the cycle a missing AMO finalizes at S1 (the S0 pulse has passed, the MSHR probe does not see the entry yet), letting a same-line request link behind it. A registered S1 allocation term closes it, in both the LLC and passthrough guards, and the MSHR asserts that no request links behind a pending AMO at the LLC. - chain_stall now paces any AMO behind a busy compute stage, not only a same-line one: a core AMO admitted while a chain's head leaves through the forward port could otherwise be overwritten by the chain's storing AMO. - A flush no longer starts while AMO results are still queued (bank_empty now includes commit_busy). SimX models no writeback queue; cycles change only when a fill would overflow it. Verified: amo and cache cells on rtlsim, XLEN 32 and 64, all pass (zacas included). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…able type
The shared Ramulator wrapper (sim/common/dram_sim) capped every simulator's
memory far below the configured device:
- one injection queue for all channels accepted one request per DRAM cycle
in total, and one response was drained per core cycle device-wide, so the
whole memory delivered at most about one line per cycle;
- each 64 B line was split into four 16 B requests, four bursts' worth of
command bandwidth for one line;
- Ramulator's AllBank refresh manager iterates over a "rank" level, which
HBM2 and GDDR6 do not have, so refresh never fired;
- the DRAM clock was a raw -DMEM_CLOCK_RATIO define (default 1, the core
clock), unrelated to the configured platform.
DramSim now models one independent controller per platform memory bank, in
simx, rtlsim, opaesim (which also gets per-bank queues) and xrtsim alike: the
caller passes the bank as the channel and the address within it; a channel
accepts one transaction per DRAM cycle and drains its own responses; a
transaction is one burst (64 B HBM2, 32 B GDDR6) with the A100 HBM interleave
or GPGPU-Sim's RTX 3070 GDDR layout and IPOLY bank hash; refresh is issued per
channel. VX_config.toml selects it:
VX_CFG_PLATFORM_MEMORY_TYPE HBM2 | HBM2_A100 | GDDR6
VX_CFG_PLATFORM_MEMORY_CLOCK_RATE DRAM clock (MHz); with PLATFORM_CLOCK_RATE
it sets the DRAM cycles per core cycle
VX_CFG_PLATFORM_MEMORY_LATENCY fixed read latency (core cycles)
VX_CFG_PLATFORM_MEMORY_WINDOW requests a channel holds in flight
The refresh manager and Ramulator's GDDR6 device (shipped but not built
upstream) are C++20 like Ramulator, so they live in a companion library,
third_party/ramulator_ext -> libramulator-vortex.so, built by
third_party/Makefile and part of the CI third-party cache key.
tests/unittest/dram_sim checks the properties the old model broke: bandwidth
scales with the channel count (40 channels = 40x one), the fixed latency adds
exactly, the window bounds requests in flight, a two-burst block answers once,
refresh stalls reads (fails with the AllBank manager), and the A100 and RTX
5080 shapes land on their targets: A100 open-row/row-conflict reads take
216/257 core cycles (GPGPU-Sim SM80: ~215/~255) and stream 952 of 1,373
B/cycle; GDDR6 at the 5080 shape streams 393 of its 418 B/cycle peak (the
card's 960 GB/s at the 2,295 MHz base clock).
Every rtlsim/simx cycle count moves; the perf baselines are re-recorded in
the same series.
Verified: tests/unittest/dram_sim passes (12 checks); every simulator builds
with the new model at XLEN 32 and 64; the CI matrix reruns on the series.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Adds the named machines next to rtx5080 (renamed from rtx_5080), each with a source for every value and sizes rounded to the powers of two the caches accept: - tinygpu: 1 core, 2 warps x 2 threads, shared memory at its default size, default ISA (so it runs the shipped RV32IMAF/RV64IMAFD runtime); the FPGA bring-up and fast simulation machine; - rtx1080 (GTX 1080), v100, a100, h100 and rtx5080 as SimX-scale study machines, now including their DRAM: HBM2 for V100, A100 HBM timing, HBM3 bandwidth on the HBM2 model for H100, GDDR5X/GDDR7 bandwidth on the GDDR6 model for the GeForce parts. Documents profiles and the DRAM keys in docs/simulation.md, configure --profile in docs/install_vortex.md and the profile layer in the build configuration design. Verified: every profile builds and passes vecadd on simx (XLEN 32); tinygpu passes sgemm on rtlsim at XLEN 32 and 64 with the same instruction count as simx, and sgemm2 (-t2) and dotproduct, which go through shared memory, on simx and rtlsim at XLEN 32 and 64. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
regression.sh fell back to $HOME/tools/gem5 when GEM5_HOME was unset, but
ci/gem5_install.sh installs gem5 under $TOOLDIR/gem5, so the gem5 cell failed
in any build configured with another --tooldir ("gem5.opt not found"). CI
sets GEM5_HOME explicitly and is unaffected. The fallback now follows
TOOLDIR, defaulting to the configure-time value.
Verified: the gem5 cell passes at XLEN 64 in a build configured with a
non-default --tooldir.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Regenerated with `pytest ci -m perf_gate --update-baselines` (XLEN 32, rtlsim) on this series, with the toolchain the CI will download (llvm-vortex a81a5957f406). All 49 entries move, because the per-channel DRAM model changes every rtlsim cycle count: - dxa/wgmma cases: 13-20% fewer cycles (independent channels, no shared injection queue); - most core, graphics and tensor cases: within a few percent; - vecadd: +21% instructions and +24%/+32% cycles (nt4/nt16) from the global index guard the kernel now needs on a rounded-up grid (one argument load and a branch in a kernel of ~two dozen instructions); - nine entries change instruction counts, from the toolchain and the test kernel changes in this series. Verified: the 49 perf_gate cases pass on rtlsim while recording. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The kernels write a*b +/- b, which clang contracts into fmadd/fmsub/fnmadd/ fnmsub (single rounding), but the host references computed the same expressions unfused (a*b rounded first). Near a = 1, a*b - b cancels, and the reference's rounding of a*b becomes hundreds of ulp of the small result, so fmsub failed whenever the configuration's element count put a sample near that point (4 cores x 8 warps x 8 threads with an L2: up to 512 ulp). The references now use fmaf, matching the instructions the kernels exercise; the kernels are unchanged. Verified: dogfood passes 22/22 on simx at that configuration (failing before) and at the default, and on rtlsim at that configuration; the kernel still emits fmadd.s x2, fmsub.s, fnmadd.s x2 and fnmsub.s. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
VX_tlb_cam and VX_mmu each defined an identical vpn_mask function. In the multi-core VM configurations (2-4 cores with an L2/L3) Verilator reported VARHIDDEN for it, which the zero-warnings policy makes fatal, so the vm:tlb-concurrent, vm:mc-l2-amo and vm:mc-l3 rtlsim builds failed. The function now lives once in VX_tlb_pkg as tlb_vpn_mask, next to tlb_perm_ok, and both modules call it. Verified: the vm cell on rtlsim passes 23/23 at XLEN 32 and 64. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The DRAM keys added with the model redesign mixed device settings with wrapper internals: a DRAM clock, a request window, an A100-specific type, per-type timing tables in DRAM cycles, a forced 256-bit channel and custom address maps. Studying another GPU meant re-deriving all of them. A channel is now described the way a data sheet describes it: VX_CFG_PLATFORM_MEMORY_TYPE (HBM2, HBM3, GDDR6), _DATA_RATE (MT/s per pin, default twice the core clock), _NUM_BANKS (channels), _LATENCY (fixed PHY/controller/interconnect cycles) and _QUEUE_SIZE (requests in flight). - One Ramulator memory system per channel: independent controllers, any channel count. - JEDEC organization with 32 B bursts (HBM: BL4 on 64-bit pseudo-channels; GDDR6: x16), array timings in nanoseconds converted by Ramulator at the data rate, bus timings in clocks. HBM array timing is JEDEC HBM2 as GPGPU-Sim's QV100 configuration lists it (tRCD/tCL/tRP 14, tRAS 33, tRC 47 ns); GDDR6 is Ramulator's own x16 preset (the GPGPU-Sim RTX 3070 values used before were half of real GDDR6). - Address map: four bursts per row visit, then the next pseudo-channel, bank group and bank, with the bank permuted by the folded row, so streams a power of two apart do not share banks; Ramulator's MOP4CLXOR mapper XORs with column bits instead and sent such streams to the same banks. - No refresh: GPU controllers hide refresh with per-bank refresh, for which Ramulator has no manager (its HBM REFsb also lacks the tRFCsb busy time); all-bank refresh cost every channel ~9% of its time and added 8.5 cycles to sgemm's mean memory latency. ChannelRefresh is removed; ramulator_ext only builds GDDR6.cpp. The default machine clocks the DRAM with the core: 14/26 cycles idle (open row / row conflict; master 13/27), one 64 B block per cycle at peak. SimX sgemm -n128 mean memory latency at NT=NW=4/8/16: 17.9/18.4/18.5 cycles (master 16.4/17.6/19.0). Profiles: v100 HBM2 1752 MT/s, a100 HBM2 3038 and h100 HBM3 6545 (32 channels carrying the cards' 40-channel bandwidth), rtx1080 GDDR6 10000, rtx5080 GDDR6 15000 with LATENCY 339 (407 less the 68 cycles JEDEC GDDR6 adds to a row conflict; derived, not re-measured). Streams reach 93-98% of peak. The dram_sim unit test checks idle latency, channel scaling, peak efficiency, strided streams and the queue bound. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
SimX picked the next warp greedy-then-oldest, with ages reset whenever a warp stopped requesting and ties going to the lowest warp id; the RTL's scoreboard arbiter (VX_scoreboard out_arb, round-robin, STICKY) keeps the current warp while it is ready and then passes to the next one in index order. On sgemm_tcu fp16 (NT=8) RTL warps run in lockstep, loading and waiting on memory together, while SimX warps drift apart (17/17/16/14 blocks per warp against 16 each) and overlap their memory waits. The gap grew with memory latency and reached 6.9% (tensor:model_parity-fp16, 5% tolerance) once the DRAM clock followed the core clock. SimX now uses a greedy round-robin arbiter matching VX_rr_arbiter with STICKY=1. model_parity-fp16: simx 39747 vs rtlsim 39221 cycles (+1.3%, was -6.9%); at 2000 MT/s +1.75% (was -3.8%); sgemm parity unchanged within 0.3%. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The bank pipeline is a timed FIFO whose delay is the configured latency, created with the default capacity of two, so a bank accepted only two requests per LATENCY cycles: 0.5 per cycle at the stock L2 latency of 4, far less with large caches. The RTL bank (PIPELINE_STAGES = LATENCY) accepts one request per cycle at any depth. Size the FIFO to the latency (at least two). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Regenerated with pytest ci -m perf_gate --update-baselines (49/49) after the DRAM model's data-sheet configuration: the default DRAM is now clocked with the core instead of at 1 GHz, with JEDEC timing, 32 B bursts and permuted bank interleave, so memory-bound cases moved. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.