From e200e0ec331042d217f1a97784c08838c9ff333c Mon Sep 17 00:00:00 2001 From: RunjiaChen Date: Wed, 19 Aug 2026 14:59:46 +0800 Subject: [PATCH] =?UTF-8?q?tests/regression:=20gbar=5Fphase=20=E2=80=94=20?= =?UTF-8?q?expose=20the=20global=20barrier's=20lost=20phase=20advance?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit vx_barrier.h documents gbarrier::arrive() as returning "phase (current generation number)" and gbarrier::wait(phase) as blocking "until generation > phase". Both statements only hold if a completed global-barrier generation actually advances the phase that arrive() reports. On RTL it never does. - RTL hw/rtl/core/VX_bar_unit.sv:165-170 handles the cluster response with phase_n = next_phase, and :186 store_write = req_valid || gbar_bus_if.rsp_valid commits it at the address latched by :238 store_waddr <= store_raddr. A response is only accepted while no barrier op is in execute (:57 gbar_rsp_ready = ~req_valid), so the previous cycle's read_addr was VX_wctl_unit.sv:177's fallback txbar_bus_if.data.addr, which is tied 'x with DXA off (VX_sfu_unit.sv:222-223). The flip therefore lands on an x-derived slot instead of the barrier's own slot. - SimX sim/simx/barrier_unit.cpp:152-166 global_resume() increments the phase of the same entry get_phase() reads. So the global barrier's own phase never moves on RTL. A SYNC global barrier still works, because its release is an unlock (VX_bar_unit.sv:167-168) and never consults the phase -- which is exactly why nothing else notices. But an async global arrive/wait pair can never be released by its own barrier, and any code treating the returned value as a generation number reads a constant. gbar_phase samples the probe barrier's phase, runs a full generation on it, rendezvouses the cores on a SECOND global barrier (a different id, hence a different gbar row and a different per-core slot), then samples again, and asserts the DOCUMENTED contract that the phase advanced. Sampling uses arrive() rather than wait() deliberately: wait() would block forever, turning a data mismatch into a timeout with no diagnostics. Results are grouped per core, since the flip can land on one core's slot by coincidence while every other core is stuck. The existing coverage cannot see it: tests/regression/async_gbarrier self-skips when num_cores < 2 (its main.cpp:80), VX_config.toml:5 defaults VX_CFG_NUM_CORES to 1, and ci/testcases/regression.yaml carries no multi-core shape for it, so the async-global path is never exercised at all. Measured (2 cores, 4 warps, 4 threads): ./ci/blackbox.sh --cores=2 --driver=simx --app=gbar_phase PASSED 0 -> 1 ./ci/blackbox.sh --cores=2 --driver=rtlsim --app=gbar_phase FAILED 0 -> 0 core phase@N phase@N+1 advanced (rtlsim) 0 0 0 NO 1 0 0 NO The test's Makefile defaults CONFIGS to -DVX_CFG_NUM_CORES=2, because the global barrier only exists when a cluster holds more than one core (VX_bar_unit.sv:44, USE_GBAR = VX_CFG_NUM_CORES > 1). This change is purely additive: no existing file is modified. Co-Authored-By: Claude Opus 5 --- tests/regression/gbar_phase/Makefile | 21 +++ tests/regression/gbar_phase/common.h | 23 ++++ tests/regression/gbar_phase/kernel.cpp | 49 +++++++ tests/regression/gbar_phase/main.cpp | 176 +++++++++++++++++++++++++ 4 files changed, 269 insertions(+) create mode 100644 tests/regression/gbar_phase/Makefile create mode 100644 tests/regression/gbar_phase/common.h create mode 100644 tests/regression/gbar_phase/kernel.cpp create mode 100644 tests/regression/gbar_phase/main.cpp diff --git a/tests/regression/gbar_phase/Makefile b/tests/regression/gbar_phase/Makefile new file mode 100644 index 0000000000..9b8ac90daa --- /dev/null +++ b/tests/regression/gbar_phase/Makefile @@ -0,0 +1,21 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +# The global barrier only exists when a cluster holds more than one core +# (VX_bar_unit's USE_GBAR = VX_CFG_NUM_CORES > 1), so default to 2 cores +# unless the caller pinned a core count already. +CONFIGS := $(if $(findstring -DVX_CFG_NUM_CORES,$(CONFIGS)),$(CONFIGS),$(CONFIGS) -DVX_CFG_NUM_CORES=2) + +PROJECT := gbar_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/gbar_phase/common.h b/tests/regression/gbar_phase/common.h new file mode 100644 index 0000000000..2368adceda --- /dev/null +++ b/tests/regression/gbar_phase/common.h @@ -0,0 +1,23 @@ +#ifndef _COMMON_H_ +#define _COMMON_H_ + +#include + +// Global-barrier ids used by the kernel. +// +// vortex::gbarrier(id) packs rs1 = (id << 8) | 0x80000000, so the hardware +// derives BOTH the cluster's gbar row (rs1[8 +: NB_BITS]) and the per-core +// barrier slot ({rs1[NW_BITS-1:0], rs1[8 +: NB_BITS]}) from `id`. Two ids are +// used so the rendezvous never shares state with the barrier under probe. +#define GBAR_PROBE_ID 1 // rs1 = 0x80000100 -> gbar row 1, per-core slot 1 +#define GBAR_SYNC_ID 2 // rs1 = 0x80000200 -> gbar row 2, per-core slot 2 + +typedef struct { + uint32_t num_cores; + uint32_t num_groups; + uint32_t group_size; + uint64_t p0_addr; // uint32_t per core: probe phase BEFORE a generation + uint64_t p1_addr; // uint32_t per core: probe phase AFTER that generation +} kernel_arg_t; + +#endif // _COMMON_H_ diff --git a/tests/regression/gbar_phase/kernel.cpp b/tests/regression/gbar_phase/kernel.cpp new file mode 100644 index 0000000000..845bab4192 --- /dev/null +++ b/tests/regression/gbar_phase/kernel.cpp @@ -0,0 +1,49 @@ +#include +#include +#include +#include "common.h" + +// vx_barrier.h documents gbarrier::arrive() as returning "phase (current +// generation number)" and gbarrier::wait(phase) as blocking "until generation +// > phase". Both statements only hold if a completed global-barrier +// generation actually ADVANCES the phase that arrive() reports. +// +// Every warp of every core arrives on the probe barrier, so its generation +// completes; the cores then rendezvous on a SECOND global barrier (a +// different id, hence a different gbar row and a different per-core slot) so +// that the probe generation is known to have finished everywhere; then the +// probe phase is sampled again. +// +// Sampling is done with arrive() rather than wait() on purpose: wait() would +// BLOCK forever if the phase never advances, turning a data mismatch into a +// timeout with no diagnostics. arrive() is non-blocking, so the test always +// terminates and always reports the two observed phases. +__kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) { + vortex::gbarrier probe(GBAR_PROBE_ID); + vortex::gbarrier sync(GBAR_SYNC_ID); + + auto p0_ptr = reinterpret_cast(arg->p0_addr); + auto p1_ptr = reinterpret_cast(arg->p1_addr); + + auto cid = vx_core_id(); + auto wid = vx_warp_id(); + auto tid = vx_thread_id(); + + // Generation N of the probe barrier: every active warp of every core + // arrives, so every core forwards and the cluster releases. + uint32_t p0 = probe.arrive(); + + // Rendezvous on the other global barrier. Its own round trip through the + // cluster unit is issued after the probe's, so once this returns the probe + // generation has completed on every core. + sync.arrive_and_wait(); + + // Sampled at the arrival of generation N+1, i.e. after generation N + // completed: the documented phase must have moved. + uint32_t p1 = probe.arrive(); + + if (wid == 0 && tid == 0) { + p0_ptr[cid] = p0; + p1_ptr[cid] = p1; + } +} diff --git a/tests/regression/gbar_phase/main.cpp b/tests/regression/gbar_phase/main.cpp new file mode 100644 index 0000000000..9a0a77adbc --- /dev/null +++ b/tests/regression/gbar_phase/main.cpp @@ -0,0 +1,176 @@ +#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 p0_buffer = nullptr; +vx_buffer_h p1_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 (p0_buffer) vx_buffer_release(p0_buffer); + if (p1_buffer) vx_buffer_release(p1_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_cores < 2) { + std::cout << "Device does not have enough cores to run the test (need at least 2)" << std::endl; + cleanup(); + return -1; + } + + kernel_arg.num_cores = static_cast(num_cores); + kernel_arg.num_groups = static_cast(num_cores * num_warps); + kernel_arg.group_size = static_cast(num_threads); + + uint32_t buf_size = kernel_arg.num_cores * sizeof(uint32_t); + + std::cout << "num_cores=" << num_cores + << ", num_warps=" << num_warps + << ", num_threads=" << num_threads << std::endl; + std::cout << "num_groups=" << kernel_arg.num_groups << std::endl; + + std::cout << "allocate device memory" << std::endl; + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ_WRITE, &p0_buffer)); + RT_CHECK(vx_buffer_address(p0_buffer, &kernel_arg.p0_addr)); + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ_WRITE, &p1_buffer)); + RT_CHECK(vx_buffer_address(p1_buffer, &kernel_arg.p1_addr)); + + // 0xff is not a legal 1-bit phase, so a slot the kernel never wrote is + // distinguishable from a phase of 0. + std::vector h_p0(kernel_arg.num_cores, 0xff); + std::vector h_p1(kernel_arg.num_cores, 0xff); + + RT_CHECK(vx_enqueue_write(queue, p0_buffer, 0, h_p0.data(), buf_size, 0, nullptr, nullptr)); + RT_CHECK(vx_enqueue_write(queue, p1_buffer, 0, h_p1.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; + { + uint32_t grid_dim[1] = {kernel_arg.num_groups}; + uint32_t block_dim[1] = {kernel_arg.group_size}; + 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] = grid_dim[0]; + li.block_dim[0] = block_dim[0]; + RT_CHECK(vx_enqueue_launch(queue, &li, 0, nullptr, &launch_ev)); + } + + std::cout << "download results" << std::endl; + vx_event_h read_ev0 = nullptr, read_ev1 = nullptr; + RT_CHECK(vx_enqueue_read(queue, h_p0.data(), p0_buffer, 0, buf_size, 1, &launch_ev, &read_ev0)); + RT_CHECK(vx_enqueue_read(queue, h_p1.data(), p1_buffer, 0, buf_size, 1, &read_ev0, &read_ev1)); + + std::cout << "wait for completion" << std::endl; + RT_CHECK(vx_event_wait_value(read_ev1, 1, VX_TIMEOUT_INFINITE)); + vx_event_release(read_ev1); + vx_event_release(read_ev0); + vx_event_release(launch_ev); + + // A global barrier is a per-CLUSTER object, so the phase is reported per + // core: a per-core table is what makes the failure legible (the flip may + // land on one core's slot by coincidence while every other core is stuck). + int errors = 0; + std::cout << "\ncore phase@N phase@N+1 advanced" << std::endl; + for (uint32_t i = 0; i < kernel_arg.num_cores; ++i) { + bool advanced = (h_p0[i] != h_p1[i]); + std::cout << " " << i + << " " << h_p0[i] + << " " << h_p1[i] + << " " << (advanced ? "yes" : "NO") << std::endl; + if (h_p0[i] == 0xff || h_p1[i] == 0xff) { + std::cout << " core " << i << ": kernel never wrote a phase" << std::endl; + ++errors; + } else if (!advanced) { + // vx_barrier.h: arrive() returns "phase (current generation number)". + // A completed generation must therefore change it. + std::cout << " core " << i << ": PHASE error: a completed global-barrier" + << " generation did not advance the phase (" << h_p0[i] + << " -> " << h_p1[i] << ")" << std::endl; + ++errors; + } + } + std::cout << "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 errors; + } + + std::cout << "PASSED!" << std::endl; + return 0; +}