fix the bug of global barrier and add the tests that can reproduce the previous bug.

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:

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",

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

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
 

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

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_