From ee0a7325c622a796663fe1a9ae48c61a8f0ebfbf Mon Sep 17 00:00:00 2001 From: RunjiaChen Date: Sun, 30 Aug 2026 11:54:27 +0800 Subject: [PATCH] =?UTF-8?q?tests/regression:=20bar=5Fslot=5Fphase=20?= =?UTF-8?q?=E2=80=94=20expose=20cross-slot=20barrier=20phase=20corruption?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A count-1 barrier arrival completes its generation immediately, so it must advance that slot's phase: the next arrival on the same slot must observe the complement. Barrier slots are independent state, so a concurrent arrival on a DIFFERENT slot must not disturb it. On RTL it does. hw/rtl/core/VX_bar_unit.sv:236 phase_r <= store_write ? store_phase_wdata : store_phase_rdata_v; reloads the working phase from whatever was just written, without checking that the write targeted the slot now being read -- unlike the rd-return path, which IS address-qualified (:224 is_rdw_hazard, :243 phase_async). The phase store is the async-read, write-first VX_dp_ram, so when two barrier requests for DIFFERENT slots are processed on consecutive cycles, the second one's phase compare, flip and write-back all use the FIRST slot's just-written phase. Its own slot is left unchanged. Requests can be back-to-back: VX_wctl_unit's wctl_reg is a plain one-per-cycle pipe with enable=1, and nothing serialises different warps' BAR-class ops. SimX keeps a phase per slot (sim/simx/barrier_unit.cpp) and is unaffected. Only the async barrier API can see this. A sync vx.bar releases on the arrival count and never compares the phase, so the corrupted bit is invisible there; it is the phase token handed out by barrier::arrive() and consumed by barrier::wait() that goes stale, which can park a warp forever or release it a generation early. bar_slot_phase releases warp 0 and warp 1 together on a gate barrier so the scheduler issues their arrivals on consecutive cycles, on two different slots, and has each warp audit its OWN slot: pre = arrive(), post = arrive(), and post must be the complement of pre. Two details make the observation stable: the two probe slots must hold EQUAL phases for the cross-slot interaction to be observable, so a detected miss is followed by one extra arrival to restore the parity for the next round; and `pre` and `post` are separated by nops so the audit read is never itself the back-to-back partner of the arrival it is auditing. The existing coverage cannot see it: tests/regression/async_barrier uses one barrier id per phase of its pipeline and never puts two different slots in flight in the same cycle window, so every barrier request it issues is either isolated or targets the slot that was just written. Measured on upstream master @ 5d62846c6 (clean tree, cold build) (2 cores, 4 warps, 4 threads, num_barriers=8, 256 rounds): ./ci/blackbox.sh --cores=2 --driver=simx --app=bar_slot_phase PASSED ./ci/blackbox.sh --cores=2 --driver=rtlsim --app=bar_slot_phase FAILED core warp missed-flips first(pre,post) (rtlsim) 0 0 0 (0,0) <- issues first, unaffected 0 1 256 (0,0) <- issues one cycle later, other slot 1 0 0 (0,0) 1 1 256 (0,0) 512 errors: every one of the 256 audited rounds on the second-issuing warp, on both cores. The stock-config simx arm must pass; that is what proves the audit itself is sound. This change is purely additive: no existing file is modified, and bar_slot_phase is deliberately NOT added to the TESTS list in tests/regression/Makefile, so CI is unaffected by a test that is expected to fail until a fix lands. Co-Authored-By: Claude Opus 5 --- tests/regression/bar_slot_phase/Makefile | 16 ++ tests/regression/bar_slot_phase/common.h | 24 +++ tests/regression/bar_slot_phase/kernel.cpp | 70 ++++++++ tests/regression/bar_slot_phase/main.cpp | 182 +++++++++++++++++++++ 4 files changed, 292 insertions(+) create mode 100644 tests/regression/bar_slot_phase/Makefile create mode 100644 tests/regression/bar_slot_phase/common.h create mode 100644 tests/regression/bar_slot_phase/kernel.cpp create mode 100644 tests/regression/bar_slot_phase/main.cpp diff --git a/tests/regression/bar_slot_phase/Makefile b/tests/regression/bar_slot_phase/Makefile new file mode 100644 index 0000000000..327948b9ce --- /dev/null +++ b/tests/regression/bar_slot_phase/Makefile @@ -0,0 +1,16 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +PROJECT := bar_slot_phase + +SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT) + +SRCS := $(SRC_DIR)/main.cpp + +VX_SRCS := $(SRC_DIR)/kernel.cpp + +OPTS ?= + +KERNEL_LIB := vortex2 + +include ../common.mk diff --git a/tests/regression/bar_slot_phase/common.h b/tests/regression/bar_slot_phase/common.h new file mode 100644 index 0000000000..b7ebaa9cd3 --- /dev/null +++ b/tests/regression/bar_slot_phase/common.h @@ -0,0 +1,24 @@ +#ifndef _COMMON_H_ +#define _COMMON_H_ + +#include + +// Barrier ids. vortex::group_barrier(id) packs rs1 = (id << 8), and the +// hardware slot address is {rs1[NW_BITS-1:0], rs1[8 +: NB_BITS]}, so these +// three ids are three DISTINCT per-core barrier slots. +#define GATE_ID 2 // rendezvous: releases every warp of the core at once +#define PROBE0_ID 0 // warp 0 probes this slot +#define PROBE1_ID 1 // warp 1 probes this slot + +#define ROUNDS 256 + +typedef struct { + uint32_t num_cores; + uint32_t num_warps; + uint32_t rounds; + uint64_t err_addr; // uint32_t per (core,warp): rounds where the phase failed to flip + uint64_t pre_addr; // uint32_t per (core,warp): first observed `pre` on failure + uint64_t post_addr; // uint32_t per (core,warp): first observed `post` on failure +} kernel_arg_t; + +#endif // _COMMON_H_ diff --git a/tests/regression/bar_slot_phase/kernel.cpp b/tests/regression/bar_slot_phase/kernel.cpp new file mode 100644 index 0000000000..26699c6f70 --- /dev/null +++ b/tests/regression/bar_slot_phase/kernel.cpp @@ -0,0 +1,70 @@ +#include +#include +#include +#include "common.h" + +// A count-1 barrier arrival completes its generation immediately, so it must +// advance that slot's phase: the phase returned by the NEXT arrival on the +// same slot must be the complement of the one this arrival returned. That is +// a per-slot invariant -- barrier slots are independent state. +// +// The invariant is probed while a SECOND warp arrives on a DIFFERENT slot in +// the same cycle window: warp 0 and warp 1 are released together by a gate +// barrier, so the scheduler issues their arrivals on consecutive cycles. +// +// Spacing: `pre` and `post` are separated by nops so the audit read is never +// itself the back-to-back partner of the arrival it is auditing. +// +// Self-correction: a detected miss leaves the slot one flip short, which +// would desynchronise the two probe slots and hide later occurrences (the +// two slots must hold EQUAL phases for the cross-slot interaction to be +// observable at all). One extra arrival restores the parity. +#define SPACER __asm__ volatile ("nop; nop; nop; nop; nop; nop; nop; nop" ::: "memory") + +__kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) { + auto err_ptr = reinterpret_cast(arg->err_addr); + auto pre_ptr = reinterpret_cast(arg->pre_addr); + auto post_ptr = reinterpret_cast(arg->post_addr); + + auto cid = vx_core_id(); + auto wid = vx_warp_id(); + auto tid = vx_thread_id(); + + vortex::barrier gate(GATE_ID); // every warp of this core + vortex::group_barrier probe0(PROBE0_ID, 1); // count-1 -> completes at once + vortex::group_barrier probe1(PROBE1_ID, 1); + + // Exactly one warp per probe slot: two warps sharing a count-1 slot would + // see each other's flips and the audit would be meaningless. + vortex::group_barrier mine = (wid == 0) ? probe0 : probe1; + + uint32_t errors = 0, first_pre = 0, first_post = 0; + + for (uint32_t r = 0; r < arg->rounds; ++r) { + // Released on the same cycle -> warp 0 and warp 1 present their arrivals + // to the barrier unit on consecutive cycles, on different slots. + gate.arrive_and_wait(); + + if (wid < 2) { + uint32_t pre = mine.arrive(); + SPACER; + uint32_t post = mine.arrive(); + + if (((pre ^ post) & 1) == 0) { + if (errors == 0) { first_pre = pre; first_post = post; } + ++errors; + mine.arrive(); // restore the slot's parity for the next round + } + } + } + + // One rendezvous before the writes so no warp races ahead into teardown. + gate.arrive_and_wait(); + + if (tid == 0) { + uint32_t idx = cid * arg->num_warps + wid; + err_ptr[idx] = errors; + pre_ptr[idx] = first_pre; + post_ptr[idx] = first_post; + } +} diff --git a/tests/regression/bar_slot_phase/main.cpp b/tests/regression/bar_slot_phase/main.cpp new file mode 100644 index 0000000000..50db7017eb --- /dev/null +++ b/tests/regression/bar_slot_phase/main.cpp @@ -0,0 +1,182 @@ +#include +#include +#include +#include +#include +#include +#include "common.h" + +#define RT_CHECK(_expr) \ + do { \ + int _ret = _expr; \ + if (0 == _ret) \ + break; \ + printf("Error: '%s' returned %d!\n", #_expr, (int)_ret); \ + cleanup(); \ + exit(-1); \ + } while (false) + +const char* kernel_file = "kernel.vxbin"; + +vx_device_h device = nullptr; +vx_buffer_h err_buffer = nullptr; +vx_buffer_h pre_buffer = nullptr; +vx_buffer_h post_buffer = nullptr; +vx_queue_h queue = nullptr; +vx_module_h module_ = nullptr; +vx_kernel_h kernel = nullptr; +kernel_arg_t kernel_arg = {}; + +static void show_usage() { + std::cout << "Vortex Test." << std::endl; + std::cout << "Usage: [-k: kernel] [-h: help]" << std::endl; +} + +static void parse_args(int argc, char** argv) { + int c; + while ((c = getopt(argc, argv, "k:h")) != -1) { + switch (c) { + case 'k': kernel_file = optarg; break; + case 'h': show_usage(); exit(0); break; + default: show_usage(); exit(-1); + } + } +} + +void cleanup() { + if (device) { + if (err_buffer) vx_buffer_release(err_buffer); + if (pre_buffer) vx_buffer_release(pre_buffer); + if (post_buffer) vx_buffer_release(post_buffer); + if (kernel) vx_kernel_release(kernel); + if (module_) vx_module_release(module_); + if (queue) vx_queue_release(queue); + vx_device_dump_perf(device, stdout); + vx_device_release(device); + } +} + +int main(int argc, char* argv[]) { + parse_args(argc, argv); + + std::cout << "open device connection" << std::endl; + RT_CHECK(vx_device_open(0, &device)); + + vx_queue_info_t qi = { sizeof(qi), nullptr, VX_QUEUE_PRIORITY_NORMAL, 0 }; + RT_CHECK(vx_queue_create(device, &qi, &queue)); + + uint64_t num_cores, num_warps, num_threads; + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_CORES, &num_cores)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_WARPS, &num_warps)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_THREADS, &num_threads)); + + if (num_warps < 2) { + std::cout << "Device does not have enough warps to run the test (need at least 2)" << std::endl; + cleanup(); + return -1; + } + + kernel_arg.num_cores = static_cast(num_cores); + kernel_arg.num_warps = static_cast(num_warps); + kernel_arg.rounds = ROUNDS; + + // One block per core, sized to the whole core, so that a block's barrier + // (vortex::barrier, default num_warps = get_num_sub_groups()) rendezvouses + // every warp of that core -- which is what releases warp 0 and warp 1 on + // the same cycle. + uint32_t num_slots = kernel_arg.num_cores * kernel_arg.num_warps; + uint32_t buf_size = num_slots * sizeof(uint32_t); + + std::cout << "num_cores=" << num_cores + << ", num_warps=" << num_warps + << ", num_threads=" << num_threads + << ", rounds=" << kernel_arg.rounds << std::endl; + + std::cout << "allocate device memory" << std::endl; + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ_WRITE, &err_buffer)); + RT_CHECK(vx_buffer_address(err_buffer, &kernel_arg.err_addr)); + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ_WRITE, &pre_buffer)); + RT_CHECK(vx_buffer_address(pre_buffer, &kernel_arg.pre_addr)); + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ_WRITE, &post_buffer)); + RT_CHECK(vx_buffer_address(post_buffer, &kernel_arg.post_addr)); + + std::vector h_err(num_slots, 0xffffffff); + std::vector h_pre(num_slots, 0xffffffff); + std::vector h_post(num_slots, 0xffffffff); + + RT_CHECK(vx_enqueue_write(queue, err_buffer, 0, h_err.data(), buf_size, 0, nullptr, nullptr)); + RT_CHECK(vx_enqueue_write(queue, pre_buffer, 0, h_pre.data(), buf_size, 0, nullptr, nullptr)); + RT_CHECK(vx_enqueue_write(queue, post_buffer, 0, h_post.data(), buf_size, 0, nullptr, nullptr)); + + std::cout << "load kernel module" << std::endl; + RT_CHECK(vx_module_load_file(device, kernel_file, &module_)); + RT_CHECK(vx_module_get_kernel(module_, "main", &kernel)); + + std::cout << "start device" << std::endl; + vx_event_h launch_ev = nullptr; + { + vx_launch_info_t li = {}; + li.struct_size = sizeof(li); + li.kernel = kernel; + li.args_host = &kernel_arg; + li.args_size = sizeof(kernel_arg); + li.ndim = 1; + li.grid_dim[0] = kernel_arg.num_cores; + li.block_dim[0] = static_cast(num_warps * num_threads); + RT_CHECK(vx_enqueue_launch(queue, &li, 0, nullptr, &launch_ev)); + } + + std::cout << "download results" << std::endl; + vx_event_h ev0 = nullptr, ev1 = nullptr, ev2 = nullptr; + RT_CHECK(vx_enqueue_read(queue, h_err.data(), err_buffer, 0, buf_size, 1, &launch_ev, &ev0)); + RT_CHECK(vx_enqueue_read(queue, h_pre.data(), pre_buffer, 0, buf_size, 1, &ev0, &ev1)); + RT_CHECK(vx_enqueue_read(queue, h_post.data(), post_buffer, 0, buf_size, 1, &ev1, &ev2)); + + std::cout << "wait for completion" << std::endl; + RT_CHECK(vx_event_wait_value(ev2, 1, VX_TIMEOUT_INFINITE)); + vx_event_release(ev2); + vx_event_release(ev1); + vx_event_release(ev0); + vx_event_release(launch_ev); + + // Only warps 0 and 1 probe; every warp writes its slot, so an unwritten + // entry means the launch did not cover the core the way the test assumes. + int errors = 0; + std::cout << "\ncore warp missed-flips first(pre,post)" << std::endl; + for (uint32_t c = 0; c < kernel_arg.num_cores; ++c) { + for (uint32_t w = 0; w < kernel_arg.num_warps; ++w) { + uint32_t i = c * kernel_arg.num_warps + w; + if (h_err[i] == 0xffffffff) { + std::cout << " " << c << " " << w << " " << std::endl; + ++errors; + continue; + } + if (w < 2) { + std::cout << " " << c << " " << w << " " << h_err[i] + << " (" << h_pre[i] << "," << h_post[i] << ")" << std::endl; + } + if (h_err[i] != 0) { + // A count-1 arrival completes its generation, so the next arrival on + // that slot must observe the complemented phase. + std::cout << " core " << c << " warp " << w << ": SLOT PHASE error: " + << h_err[i] << " of " << kernel_arg.rounds + << " rounds saw a count-1 arrival fail to advance its own slot's" + << " phase (pre=" << h_pre[i] << " post=" << h_post[i] << ")" << std::endl; + errors += static_cast(h_err[i]); + } + } + } + std::cout << "slot phase errors: " << errors << std::endl; + + std::cout << "cleanup" << std::endl; + cleanup(); + + if (errors != 0) { + std::cout << "Found " << errors << " errors!" << std::endl; + std::cout << "FAILED!" << std::endl; + return 1; + } + + std::cout << "PASSED!" << std::endl; + return 0; +}