diff --git a/ci/testcases/regression.yaml b/ci/testcases/regression.yaml index 94ca0e12c7..3250b9fe88 100644 --- a/ci/testcases/regression.yaml +++ b/ci/testcases/regression.yaml @@ -76,6 +76,20 @@ 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: + - rtlsim + app: gbar_phase + shape: + cores: 2 - id: vecadd-1 via: blackbox drivers: 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; } } } diff --git a/tests/regression/Makefile b/tests/regression/Makefile index cccdee5116..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 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/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); } 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; +}