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; +}