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