From 451dded2af7b4ff584be677b3324448e1bc763ff Mon Sep 17 00:00:00 2001 From: symmetryyyyy Date: Wed, 2 Sep 2026 16:09:44 -0700 Subject: [PATCH 1/4] tests: cover barrier phase routing Add non-deadlocking diagnostics for cross-slot local phase corruption and missing global-barrier phase advancement. Register both tests in the aggregate regression suite and explicit SimX/xrtsim smoke coverage. Co-Authored-By: RunjiaChen Co-Authored-By: Claude Opus 5 --- ci/testcases/regression.yaml | 28 ++++ tests/regression/Makefile | 2 +- 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 +++++++++++++++++++++ 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 ++++++++++++++++++++ 10 files changed, 590 insertions(+), 1 deletion(-) 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 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/ci/testcases/regression.yaml b/ci/testcases/regression.yaml index 94ca0e12c7..a9cfbe1cca 100644 --- a/ci/testcases/regression.yaml +++ b/ci/testcases/regression.yaml @@ -76,6 +76,34 @@ tests: - xrt app: dogfood args: -n1 -tbar +- id: gbar-phase-1 + via: blackbox + drivers: + - simx + app: gbar_phase + shape: + cores: 2 +- id: gbar-phase-2 + via: blackbox + drivers: + - xrt + app: gbar_phase + shape: + cores: 2 +- id: bar-slot-phase-1 + via: blackbox + drivers: + - simx + app: bar_slot_phase + shape: + warps: 2 +- id: bar-slot-phase-2 + via: blackbox + drivers: + - xrt + app: bar_slot_phase + shape: + warps: 2 - id: vecadd-1 via: blackbox drivers: diff --git a/tests/regression/Makefile b/tests/regression/Makefile index cccdee5116..828c65a6f7 100644 --- a/tests/regression/Makefile +++ b/tests/regression/Makefile @@ -13,7 +13,7 @@ TESTS := \ sgemm2_dxa sgemm2_tcu sgemm_tcu_wg_dxa sgemm_tcu_wg_sp_dxa \ sgemm2_dxa_mcast sgemm_tcu_wg_dxa_mcast \ dxa_copy dxa_copy_mcast dxa_kmajor_check \ - async_barrier async_gbarrier packld wgather wsync quad_lod \ + async_barrier async_gbarrier gbar_phase bar_slot_phase packld wgather wsync quad_lod \ occupancy multikernel reopen module_reload single_cta \ vm_test vm_fault 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; +} 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; +} From 3b5b6f7bc842db88faa47f82d8f2c8317863406d Mon Sep 17 00:00:00 2001 From: symmetryyyyy Date: Wed, 2 Sep 2026 16:05:59 -0700 Subject: [PATCH 2/4] tests: widen the asynchronous global-barrier window --- tests/regression/async_gbarrier/kernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/regression/async_gbarrier/kernel.cpp b/tests/regression/async_gbarrier/kernel.cpp index 7f9f435f58..af0c478b92 100644 --- a/tests/regression/async_gbarrier/kernel.cpp +++ b/tests/regression/async_gbarrier/kernel.cpp @@ -25,7 +25,7 @@ __kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) { uint32_t phase = bar1.arrive(); uint32_t overlap_work = ((cid + 1) * 0x9e3779b9u) ^ ((wid + 1) * 0x85ebca6bu) ^ (tid + 1); - for (uint32_t i = 0; i < 32; ++i) { + for (uint32_t i = 0; i < 1024; ++i) { overlap_work = overlap_work * 1664525u + 1013904223u; overlap_work ^= (overlap_work >> 13); } From 9e5f577e20e2ecc62c979ced95350b4e13a7ab27 Mon Sep 17 00:00:00 2001 From: symmetryyyyy Date: Wed, 2 Sep 2026 16:06:03 -0700 Subject: [PATCH 3/4] fix(gbar): retire responses by barrier id Capture the phase and exact waiters for each pending global-barrier generation, queue completed local arrivals, and route each cluster response directly to its barrier slot. Keep response-only writes out of the barrier state RAM and qualify the working phase bypass by address. Use one-bit phase toggles in SimX so wait comparisons keep matching the RTL after wraparound. --- hw/rtl/core/VX_bar_unit.sv | 139 ++++++++++++++++++++++++------------- sim/simx/barrier_unit.cpp | 6 +- 2 files changed, 92 insertions(+), 53 deletions(-) diff --git a/hw/rtl/core/VX_bar_unit.sv b/hw/rtl/core/VX_bar_unit.sv index eaebb44497..bf7989b166 100644 --- a/hw/rtl/core/VX_bar_unit.sv +++ b/hw/rtl/core/VX_bar_unit.sv @@ -41,6 +41,7 @@ module VX_bar_unit import VX_gpu_pkg::*; #( // warp mask + warp count + event count localparam EVENT_WIDTH = `CLOG2(`VX_CFG_MAX_BAR_EVENTS + 1); localparam BAR_STATEW = `VX_CFG_NUM_WARPS + NW_WIDTH + EVENT_WIDTH; + localparam GBAR_REQW = NB_WIDTH + NC_WIDTH; localparam USE_GBAR = (`VX_CFG_NUM_CORES > 1); logic [`VX_CFG_NUM_WARPS-1:0] mask_r, mask_n; @@ -51,10 +52,18 @@ module VX_bar_unit import VX_gpu_pkg::*; #( logic unlock_valid_n; logic [`VX_CFG_NUM_WARPS-1:0] unlock_mask_n; - logic gbar_req_valid_r, gbar_req_valid_n; - logic [NB_WIDTH-1:0] gbar_req_id_r, gbar_req_id_n; - logic [NC_WIDTH-1:0] gbar_req_size_m1_r, gbar_req_size_m1_n; + // A response carries only its id after the barrier-store read port has moved on. + logic [`VX_CFG_NUM_BARRIERS-1:0] gbar_pending_r, gbar_pending_n; + logic [`VX_CFG_NUM_BARRIERS-1:0] gbar_pending_phase_r, gbar_pending_phase_n; + logic [`VX_CFG_NUM_BARRIERS-1:0][`VX_CFG_NUM_WARPS-1:0] gbar_waiters_r, gbar_waiters_n; + + logic gbar_enqueue; + logic [NB_WIDTH-1:0] gbar_enqueue_id; + logic [NC_WIDTH-1:0] gbar_enqueue_size_m1; + wire gbar_rsp_ready = ~req_valid; + wire gbar_rsp_fire = gbar_bus_if.rsp_valid && gbar_rsp_ready; + wire gbar_rsp_apply = USE_GBAR && gbar_rsp_fire && gbar_pending_r[gbar_bus_if.rsp_data.id]; wire [`VX_CFG_NUM_WARPS-1:0] wait_mask = ((`VX_CFG_NUM_WARPS)'(1) << req_wid) | mask_r; wire [NW_WIDTH-1:0] next_count = count_r + NW_WIDTH'(1); @@ -67,9 +76,12 @@ module VX_bar_unit import VX_gpu_pkg::*; #( phase_n = phase_r; unlock_valid_n = 0; unlock_mask_n = 'x; - gbar_req_valid_n = gbar_req_valid_r; - gbar_req_id_n = gbar_req_id_r; - gbar_req_size_m1_n = gbar_req_size_m1_r; + gbar_pending_n = gbar_pending_r; + gbar_pending_phase_n = gbar_pending_phase_r; + gbar_waiters_n = gbar_waiters_r; + gbar_enqueue = 0; + gbar_enqueue_id = 'x; + gbar_enqueue_size_m1 = 'x; // local barrier scheduling if (req_valid && ~req_data.is_global) begin @@ -134,20 +146,23 @@ module VX_bar_unit import VX_gpu_pkg::*; #( // unlock warps if decrementing event to 0 and all warps have arrived if ((req_data.phase == 0) && (events_r == EVENT_WIDTH'(1)) && (wait_mask == active_warps)) begin mask_n = '0; - gbar_req_valid_n = 1; // notify global barrier - gbar_req_id_n = req_data.id; - gbar_req_size_m1_n = NC_WIDTH'(count_r); // was saved in barrier_arrive + gbar_enqueue = 1; // notify global barrier + gbar_enqueue_id = req_data.id; + gbar_enqueue_size_m1 = NC_WIDTH'(count_r); // was saved in barrier_arrive end end else if (req_data.is_arrive) begin // barrier arrival count_n = NW_WIDTH'(req_data.size_m1); // store participating number of cores + if (req_data.is_sync) begin + gbar_waiters_n[req_data.id][req_wid] = 1; + end if (wait_mask == active_warps && events_r == 0) begin mask_n = '0; - gbar_req_valid_n = 1; // notify global barrier - gbar_req_id_n = req_data.id; - gbar_req_size_m1_n = NC_WIDTH'(req_data.size_m1); + gbar_enqueue = 1; // notify global barrier + gbar_enqueue_id = req_data.id; + gbar_enqueue_size_m1 = NC_WIDTH'(req_data.size_m1); end else begin - // Add arriving warp to wait mask + // Add arriving warp to arrival mask mask_n = wait_mask; end end else begin @@ -157,21 +172,22 @@ module VX_bar_unit import VX_gpu_pkg::*; #( unlock_mask_n = (`VX_CFG_NUM_WARPS)'(1) << req_wid; end else begin // add warp to wait list - mask_n = wait_mask; + gbar_waiters_n[req_data.id][req_wid] = 1; end end end // global barrier response handling - if (gbar_bus_if.rsp_valid && gbar_rsp_ready && (gbar_bus_if.rsp_data.id == gbar_req_id_r)) begin - unlock_valid_n = 1; // release stalled warps - unlock_mask_n = active_warps; // release all active warps - phase_n = next_phase; // advance phase + if (gbar_rsp_apply) begin + unlock_valid_n = (gbar_waiters_r[gbar_bus_if.rsp_data.id] != '0); + unlock_mask_n = gbar_waiters_r[gbar_bus_if.rsp_data.id]; + gbar_waiters_n[gbar_bus_if.rsp_data.id] = '0; + gbar_pending_n[gbar_bus_if.rsp_data.id] = 0; end - // global barrier request handshake - if (gbar_req_valid_r && gbar_bus_if.req_ready) begin - gbar_req_valid_n = 0; + if (gbar_enqueue) begin + gbar_pending_n[gbar_enqueue_id] = 1; + gbar_pending_phase_n[gbar_enqueue_id] = phase_r; end end end @@ -182,8 +198,10 @@ module VX_bar_unit import VX_gpu_pkg::*; #( wire [BAR_ADDR_W-1:0] store_raddr = read_addr; reg [BAR_ADDR_W-1:0] store_waddr; wire [BAR_STATEW-1:0] store_state_wdata = {mask_n, count_n, events_n}; - wire store_phase_wdata = phase_n; - wire store_write = req_valid || gbar_bus_if.rsp_valid; + wire store_state_write = req_valid; + wire [BAR_ADDR_W-1:0] store_phase_waddr = gbar_rsp_apply ? BAR_ADDR_W'(gbar_bus_if.rsp_data.id) : store_waddr; + wire store_phase_wdata = gbar_rsp_apply ? ~gbar_pending_phase_r[gbar_bus_if.rsp_data.id] : phase_n; + wire store_phase_write = req_valid || gbar_rsp_apply; VX_dp_ram #( .DATAW (BAR_STATEW), @@ -194,7 +212,7 @@ module VX_bar_unit import VX_gpu_pkg::*; #( .clk (clk), .reset (reset), .read (1'b1), - .write (store_write), + .write (store_state_write), .wren (1'b1), .raddr (store_raddr), .waddr (store_waddr), @@ -211,17 +229,17 @@ module VX_bar_unit import VX_gpu_pkg::*; #( .clk (clk), .reset (reset), .read (1'b1), - .write (store_write), + .write (store_phase_write), .wren (1'b1), .raddr (store_raddr), - .waddr (store_waddr), + .waddr (store_phase_waddr), .wdata (store_phase_wdata), .rdata (store_phase_rdata) ); // Store reset handling reg [(1 << BAR_ADDR_BITS)-1:0] store_valids; - wire is_rdw_hazard = store_write && (store_waddr == store_raddr); + wire is_phase_rdw_hazard = store_phase_write && (store_phase_waddr == store_raddr); wire store_phase_rdata_v = store_valids[store_raddr] ? store_phase_rdata : '0; @@ -230,17 +248,17 @@ module VX_bar_unit import VX_gpu_pkg::*; #( store_valids <= '0; phase_r <= '0; end else begin - if (store_write) begin + if (store_state_write) begin store_valids[store_waddr] <= 1'b1; end - phase_r <= store_write ? store_phase_wdata : store_phase_rdata_v; + phase_r <= is_phase_rdw_hazard ? store_phase_wdata : store_phase_rdata_v; end store_waddr <= store_raddr; end assign {mask_r, count_r, events_r} = store_valids[store_waddr] ? store_state_rdata : '0; - wire phase_async = is_rdw_hazard ? phase_n : store_phase_rdata_v; + wire phase_async = is_phase_rdw_hazard ? store_phase_wdata : store_phase_rdata_v; reg unlock_valid_r; reg [`VX_CFG_NUM_WARPS-1:0] unlock_mask_r; @@ -258,35 +276,56 @@ module VX_bar_unit import VX_gpu_pkg::*; #( assign unlock_valid = unlock_valid_r; assign unlock_mask = unlock_mask_r; - if (USE_GBAR) begin : g_gbar - - always @(posedge clk) begin - if (reset) begin - gbar_req_valid_r <= 0; - end else begin - gbar_req_valid_r <= gbar_req_valid_n; - end - gbar_req_size_m1_r <= gbar_req_size_m1_n; - gbar_req_id_r <= gbar_req_id_n; + always @(posedge clk) begin + if (reset) begin + gbar_pending_r <= '0; + gbar_pending_phase_r <= '0; + gbar_waiters_r <= '0; + end else begin + gbar_pending_r <= gbar_pending_n; + gbar_pending_phase_r <= gbar_pending_phase_n; + gbar_waiters_r <= gbar_waiters_n; end + end + + if (USE_GBAR) begin : g_gbar - assign gbar_bus_if.req_valid = gbar_req_valid_r; - assign gbar_bus_if.req_data.id = gbar_req_id_r; - assign gbar_bus_if.req_data.size_m1 = gbar_req_size_m1_r; + wire [GBAR_REQW-1:0] req_queue_data; + wire req_empty; + wire req_pop = ~req_empty && gbar_bus_if.req_ready; + + VX_fifo_queue #( + .DATAW (GBAR_REQW), + .DEPTH (1 << NB_WIDTH), + .LUTRAM (1) + ) req_queue ( + .clk (clk), + .reset (reset), + .push (gbar_enqueue), + .pop (req_pop), + .data_in ({gbar_enqueue_id, gbar_enqueue_size_m1}), + .data_out(req_queue_data), + .empty (req_empty), + `UNUSED_PIN (alm_empty), + `UNUSED_PIN (alm_full), + `UNUSED_PIN (full), + `UNUSED_PIN (size) + ); + + assign gbar_bus_if.req_valid = ~req_empty; + assign {gbar_bus_if.req_data.id, gbar_bus_if.req_data.size_m1} = req_queue_data; assign gbar_bus_if.req_data.core_id = NC_WIDTH'(CORE_ID % `VX_CFG_NUM_CORES); assign gbar_bus_if.rsp_ready = gbar_rsp_ready; + `RUNTIME_ASSERT(~gbar_enqueue || ~gbar_pending_r[gbar_enqueue_id], ("%s duplicate global barrier generation: id=%0d", INSTANCE_ID, gbar_enqueue_id)) + `RUNTIME_ASSERT(~gbar_enqueue || (store_waddr == BAR_ADDR_W'(gbar_enqueue_id)), ("%s invalid global barrier slot: id=%0d, slot=%0d", INSTANCE_ID, gbar_enqueue_id, store_waddr)) end else begin : g_nogbar - assign gbar_req_valid_r = 0; - assign gbar_req_size_m1_r = 'x; - assign gbar_req_id_r = 'x; - assign gbar_bus_if.req_valid = 0; assign gbar_bus_if.req_data = 'x; assign gbar_bus_if.rsp_ready = 0; - `UNUSED_VAR ({gbar_req_valid_n, gbar_req_size_m1_n, gbar_req_id_n}) + `UNUSED_VAR (gbar_enqueue_size_m1) end @@ -296,9 +335,9 @@ module VX_bar_unit import VX_gpu_pkg::*; #( `TRACE(2, ("%t: %s req: wid=%0d, bar_id=%0d, is_global=%b, is_event=%b, is_arrive=%b, is_sync=%b, phase=%b, size_m1=%0d\n", $time, INSTANCE_ID, req_wid, req_data.id, req_data.is_global, req_data.is_event, req_data.is_arrive, req_data.is_sync, req_data.phase, req_data.size_m1)) end - if (USE_GBAR && gbar_req_valid_n && ~gbar_req_valid_r) begin + if (USE_GBAR && gbar_bus_if.req_valid && gbar_bus_if.req_ready) begin `TRACE(2, ("%t: %s global-req: bar_id=%0d, size_m1=%0d\n", - $time, INSTANCE_ID, gbar_req_id_n, gbar_req_size_m1_n)) + $time, INSTANCE_ID, gbar_bus_if.req_data.id, gbar_bus_if.req_data.size_m1)) end if (USE_GBAR && gbar_bus_if.rsp_valid && gbar_rsp_ready) begin `TRACE(2, ("%t: %s global-rsp: bar_id=%0d\n", diff --git a/sim/simx/barrier_unit.cpp b/sim/simx/barrier_unit.cpp index f794e82ee4..88afeba4a7 100644 --- a/sim/simx/barrier_unit.cpp +++ b/sim/simx/barrier_unit.cpp @@ -82,7 +82,7 @@ void BarrierUnit::arrive(uint32_t bar_id, uint32_t count, uint32_t wid, bool is_ } // reset barrier and advance phase barrier.wait_mask.reset(); - ++barrier.phase; + barrier.phase ^= 1u; } // update count and wrap around if (count == 0) { @@ -123,7 +123,7 @@ void BarrierUnit::global_resume(uint32_t bar_id) { } } barrier.wait_mask.reset(); - ++barrier.phase; + barrier.phase ^= 1u; } void BarrierUnit::event_attach(uint32_t bar_id, uint32_t count) { @@ -159,7 +159,7 @@ void BarrierUnit::event_release(uint32_t bar_id) { } // reset barrier and advance phase barrier.wait_mask.reset(); - ++barrier.phase; + barrier.phase ^= 1u; } } } From bc90d020c70978b02e26656e28bef00a1453e4eb Mon Sep 17 00:00:00 2001 From: symmetryyyyy Date: Fri, 4 Sep 2026 03:43:52 -0700 Subject: [PATCH 4/4] tests: finalize global barrier regression matrix --- ci/testcases/regression.yaml | 16 +- tests/regression/Makefile | 2 +- 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 --------------------- 6 files changed, 2 insertions(+), 308 deletions(-) delete mode 100644 tests/regression/bar_slot_phase/Makefile delete mode 100644 tests/regression/bar_slot_phase/common.h delete mode 100644 tests/regression/bar_slot_phase/kernel.cpp delete mode 100644 tests/regression/bar_slot_phase/main.cpp diff --git a/ci/testcases/regression.yaml b/ci/testcases/regression.yaml index a9cfbe1cca..3250b9fe88 100644 --- a/ci/testcases/regression.yaml +++ b/ci/testcases/regression.yaml @@ -86,24 +86,10 @@ tests: - id: gbar-phase-2 via: blackbox drivers: - - xrt + - rtlsim app: gbar_phase shape: cores: 2 -- id: bar-slot-phase-1 - via: blackbox - drivers: - - simx - app: bar_slot_phase - shape: - warps: 2 -- id: bar-slot-phase-2 - via: blackbox - drivers: - - xrt - app: bar_slot_phase - shape: - warps: 2 - id: vecadd-1 via: blackbox drivers: diff --git a/tests/regression/Makefile b/tests/regression/Makefile index 828c65a6f7..103f166681 100644 --- a/tests/regression/Makefile +++ b/tests/regression/Makefile @@ -13,7 +13,7 @@ TESTS := \ sgemm2_dxa sgemm2_tcu sgemm_tcu_wg_dxa sgemm_tcu_wg_sp_dxa \ sgemm2_dxa_mcast sgemm_tcu_wg_dxa_mcast \ dxa_copy dxa_copy_mcast dxa_kmajor_check \ - async_barrier async_gbarrier gbar_phase bar_slot_phase packld wgather wsync quad_lod \ + async_barrier async_gbarrier gbar_phase packld wgather wsync quad_lod \ occupancy multikernel reopen module_reload single_cta \ vm_test vm_fault diff --git a/tests/regression/bar_slot_phase/Makefile b/tests/regression/bar_slot_phase/Makefile deleted file mode 100644 index 327948b9ce..0000000000 --- a/tests/regression/bar_slot_phase/Makefile +++ /dev/null @@ -1,16 +0,0 @@ -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 deleted file mode 100644 index b7ebaa9cd3..0000000000 --- a/tests/regression/bar_slot_phase/common.h +++ /dev/null @@ -1,24 +0,0 @@ -#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 deleted file mode 100644 index 26699c6f70..0000000000 --- a/tests/regression/bar_slot_phase/kernel.cpp +++ /dev/null @@ -1,70 +0,0 @@ -#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 deleted file mode 100644 index 50db7017eb..0000000000 --- a/tests/regression/bar_slot_phase/main.cpp +++ /dev/null @@ -1,182 +0,0 @@ -#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; -}