Bug/gbar response phase

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

tests/regression/gbar_phase/common.h

@@ -0,0 +1,23 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#include <stdint.h>
+
+// 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_

tests/regression/gbar_phase/kernel.cpp

@@ -0,0 +1,49 @@
+#include <vx_spawn2.h>
+#include <vx_intrinsics.h>
+#include <vx_barrier.h>
+#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<uint32_t*>(arg->p0_addr);
+  auto p1_ptr = reinterpret_cast<uint32_t*>(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;
+  }
+}

tests/regression/gbar_phase/main.cpp

@@ -0,0 +1,176 @@
+#include <iostream>
+#include <cstdio>
+#include <cstdlib>
+#include <unistd.h>
+#include <vector>
+#include <vortex2.h>
+#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<uint32_t>(num_cores);
+  kernel_arg.num_groups = static_cast<uint32_t>(num_cores * num_warps);
+  kernel_arg.group_size = static_cast<uint32_t>(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<uint32_t> h_p0(kernel_arg.num_cores, 0xff);
+  std::vector<uint32_t> 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;
+}