tests/regression: bar_slot_phase — expose cross-slot barrier phase co…

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

tests/regression/bar_slot_phase/common.h

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

tests/regression/bar_slot_phase/kernel.cpp

@@ -0,0 +1,70 @@
+#include <vx_spawn2.h>
+#include <vx_intrinsics.h>
+#include <vx_barrier.h>
+#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<uint32_t*>(arg->err_addr);
+  auto pre_ptr  = reinterpret_cast<uint32_t*>(arg->pre_addr);
+  auto post_ptr = reinterpret_cast<uint32_t*>(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;
+  }
+}

tests/regression/bar_slot_phase/main.cpp

@@ -0,0 +1,182 @@
+#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 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<uint32_t>(num_cores);
+  kernel_arg.num_warps = static_cast<uint32_t>(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<uint32_t> h_err(num_slots, 0xffffffff);
+  std::vector<uint32_t> h_pre(num_slots, 0xffffffff);
+  std::vector<uint32_t> 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<uint32_t>(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 << "     <not written>" << 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<int>(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;
+}