bug fixes

hw/rtl/core/VX_lsu_slice.sv

@@ -105,6 +105,8 @@ module VX_lsu_slice import VX_gpu_pkg::*; #(
     wire mem_rsp_sop_pkt, mem_rsp_eop_pkt;
     wire no_rsp_buf_valid, no_rsp_buf_ready;
 
+    wire [LSUQ_SIZEW-1:0] pkt_waddr, pkt_raddr;
+
     // fence handling
 
     reg fence_lock;
@@ -208,54 +210,33 @@ module VX_lsu_slice import VX_gpu_pkg::*; #(
         end
     end
 
-    // track SOP/EOP for out-of-order memory responses
-
-    wire [LSUQ_SIZEW-1:0] pkt_waddr, pkt_raddr;
+    // multi-packet load responses could return out-of-order.
+    // we should track and return eop packet response last.
 
     if (PID_BITS != 0) begin : g_pids
         reg [`LSUQ_IN_SIZE-1:0][PID_BITS:0] pkt_ctr;
-        reg [`LSUQ_IN_SIZE-1:0] pkt_sop, pkt_eop;
+        reg [`LSUQ_IN_SIZE-1:0] pkt_sop;
+        reg [`LSUQ_IN_SIZE-1:0] pkt_eop;
 
-        wire mem_req_rd_fire     = mem_req_fire && ~mem_req_rw;
-        wire mem_req_rd_sop_fire = mem_req_rd_fire && execute_if.data.sop;
+        wire mem_req_rd_fire = mem_req_fire && ~mem_req_rw;
         wire mem_req_rd_eop_fire = mem_req_rd_fire && execute_if.data.eop;
-        wire mem_rsp_eop_fire    = mem_rsp_fire && mem_rsp_eop;
-        wire full;
+        wire mem_rsp_eop_fire = mem_rsp_fire && mem_rsp_eop;
 
-        VX_allocator #(
-            .SIZE (`LSUQ_IN_SIZE)
-        ) pkt_allocator (
-            .clk        (clk),
-            .reset      (reset),
-            .acquire_en (mem_req_rd_eop_fire),
-            .acquire_addr(pkt_waddr),
-            .release_en (mem_rsp_eop_pkt),
-            .release_addr(pkt_raddr),
-            `UNUSED_PIN (empty),
-            .full       (full)
-        );
-
-        wire rd_during_wr = mem_req_rd_fire && mem_rsp_eop_fire && (pkt_raddr == pkt_waddr);
+        assign mem_rsp_sop_pkt = pkt_sop[pkt_raddr];
+        assign mem_rsp_eop_pkt = mem_rsp_eop && pkt_eop[pkt_raddr] && (pkt_ctr[pkt_raddr] == 1);
 
         always @(posedge clk) begin
             if (reset) begin
-                pkt_ctr <= '0;
-                pkt_sop <= '0;
-                pkt_eop <= '0;
-            end else begin
-                if (mem_req_rd_sop_fire) begin
-                    pkt_sop[pkt_waddr] <= 1;
+                for (integer i = 0; i < `LSUQ_IN_SIZE; ++i) begin
+                    pkt_ctr[i] <= '0;
+                    pkt_sop[i] <= 1;
+                    pkt_eop[i] <= 0;
                 end
+            end else begin
                 if (mem_req_rd_eop_fire) begin
                     pkt_eop[pkt_waddr] <= 1;
                 end
-                if (mem_rsp_fire) begin
-                    pkt_sop[pkt_raddr] <= 0;
-                end
-                if (mem_rsp_eop_pkt) begin
-                    pkt_eop[pkt_raddr] <= 0;
-                end
-                if (~rd_during_wr) begin
+                if (~(mem_req_rd_fire && mem_rsp_eop_fire && (pkt_raddr == pkt_waddr))) begin
                     if (mem_req_rd_fire) begin
                         pkt_ctr[pkt_waddr] <= pkt_ctr[pkt_waddr] + PID_BITS'(1);
                     end
@@ -263,16 +244,19 @@ module VX_lsu_slice import VX_gpu_pkg::*; #(
                         pkt_ctr[pkt_raddr] <= pkt_ctr[pkt_raddr] - PID_BITS'(1);
                     end
                 end
+                if (mem_rsp_fire) begin
+                    pkt_sop[pkt_raddr] <= 0;
+                end
+                if (mem_rsp_eop_fire && mem_rsp_eop_pkt) begin
+                    pkt_sop[pkt_raddr] <= 1;
+                    pkt_eop[pkt_raddr] <= 0;
+                end
             end
         end
-
-        assign mem_rsp_sop_pkt = pkt_sop[pkt_raddr];
-        assign mem_rsp_eop_pkt = mem_rsp_eop_fire && pkt_eop[pkt_raddr] && (pkt_ctr[pkt_raddr] == 1);
-        `RUNTIME_ASSERT(~(mem_req_rd_fire && full), ("%t: allocator full!", $time))
-        `RUNTIME_ASSERT(~mem_req_rd_sop_fire || 0 == pkt_ctr[pkt_waddr], ("%t: oops! broken sop request!", $time))
-        `UNUSED_VAR (mem_rsp_sop)
+        `RUNTIME_ASSERT(~(mem_req_rd_fire && pkt_eop[pkt_waddr]), ("%t: oops! broken eop request! (#%0d)", $time, execute_if.data.uuid))
+        `RUNTIME_ASSERT(~(mem_req_rd_fire && (2**PID_BITS-1) == pkt_ctr[pkt_waddr]), ("%t: oops! broken ctr request! (#%0d)", $time, execute_if.data.uuid))
+        `RUNTIME_ASSERT(~(mem_rsp_fire && 0 == pkt_ctr[pkt_raddr]), ("%t: oops! broken ctr response! (#%0d)", $time, rsp_uuid))
     end else begin : g_no_pids
-        assign pkt_waddr = 0;
         assign mem_rsp_sop_pkt = mem_rsp_sop;
         assign mem_rsp_eop_pkt = mem_rsp_eop;
         `UNUSED_VAR (pkt_raddr)
@@ -337,8 +321,12 @@ module VX_lsu_slice import VX_gpu_pkg::*; #(
         .core_req_data  (mem_req_data),
         .core_req_tag   (mem_req_tag),
         .core_req_ready (mem_req_ready),
-        `UNUSED_PIN (core_req_empty),
-        `UNUSED_PIN (core_req_wr_notify),
+        .core_req_queue_id (pkt_waddr),
+
+        // request queue info
+        `UNUSED_PIN (req_queue_empty),
+        `UNUSED_PIN (req_queue_pop),
+        `UNUSED_PIN (req_queue_id),
 
         // Output response
         .core_rsp_valid (mem_rsp_valid),

hw/rtl/libs/VX_mem_scheduler.sv

@@ -39,7 +39,8 @@ module VX_mem_scheduler #(
     parameter MEM_BATCH_BITS= `CLOG2(MEM_BATCHES),
     parameter MEM_QUEUE_ADDRW= `CLOG2(COALESCE_ENABLE ? MEM_QUEUE_SIZE : CORE_QUEUE_SIZE),
     parameter MEM_ADDR_WIDTH= ADDR_WIDTH - `CLOG2(PER_LINE_REQS),
-    parameter MEM_TAG_WIDTH = UUID_WIDTH + MEM_QUEUE_ADDRW + MEM_BATCH_BITS
+    parameter MEM_TAG_WIDTH = UUID_WIDTH + MEM_QUEUE_ADDRW + MEM_BATCH_BITS,
+    parameter CORE_QUEUE_ADDRW = `CLOG2(CORE_QUEUE_SIZE)
 ) (
     input wire clk,
     input wire reset,
@@ -54,8 +55,12 @@ module VX_mem_scheduler #(
     input wire [CORE_REQS-1:0][WORD_WIDTH-1:0] core_req_data,
     input wire [TAG_WIDTH-1:0]              core_req_tag,
     output wire                             core_req_ready,
-    output wire                             core_req_empty,
-    output wire                             core_req_wr_notify,
+    output wire [CORE_QUEUE_ADDRW-1:0]      core_req_queue_id,
+
+    // Core request queue
+    output wire                             req_queue_empty,
+    output wire                             req_queue_pop,
+    output wire [CORE_QUEUE_ADDRW-1:0]      req_queue_id,
 
     // Core response
     output wire                             core_rsp_valid,
@@ -86,7 +91,6 @@ module VX_mem_scheduler #(
 );
     localparam BATCH_SEL_WIDTH = `UP(MEM_BATCH_BITS);
     localparam STALL_TIMEOUT   = 10000000;
-    localparam CORE_QUEUE_ADDRW= `CLOG2(CORE_QUEUE_SIZE);
     localparam TAG_ID_WIDTH    = TAG_WIDTH - UUID_WIDTH;
     localparam REQQ_TAG_WIDTH  = UUID_WIDTH + CORE_QUEUE_ADDRW;
     localparam MERGED_TAG_WIDTH= UUID_WIDTH + MEM_QUEUE_ADDRW;
@@ -185,11 +189,13 @@ module VX_mem_scheduler #(
     // can accept another request?
     assign core_req_ready = reqq_ready_in && ibuf_ready;
 
-    // no pending requests
-    assign core_req_empty = !reqq_valid && ibuf_empty;
+    // return core queue id
+    assign core_req_queue_id = ibuf_waddr;
 
-    // notify write request submisison
-    assign core_req_wr_notify = reqq_valid && reqq_ready && reqq_rw;
+    // request qeueue info
+    assign req_queue_pop   = reqq_valid && reqq_ready;
+    assign req_queue_empty = !reqq_valid && ibuf_empty;
+    assign req_queue_id    = reqq_tag[CORE_QUEUE_ADDRW-1:0];
 
     // Index buffer ///////////////////////////////////////////////////////////
 

tests/regression/sgemm_tpu/Makefile

@@ -0,0 +1,14 @@
+ROOT_DIR := $(realpath ../../..)
+include $(ROOT_DIR)/config.mk
+
+PROJECT := sgemm_tpu
+
+SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT)
+
+SRCS := $(SRC_DIR)/main.cpp
+
+VX_SRCS := $(SRC_DIR)/kernel.cpp
+
+OPTS ?= -n32
+
+include ../common.mk

tests/regression/sgemm_tpu/common.h

@@ -0,0 +1,16 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#ifndef TYPE
+#define TYPE float
+#endif
+
+typedef struct {
+  uint32_t grid_dim[2];
+  uint32_t size;
+  uint64_t A_addr;
+  uint64_t B_addr;
+  uint64_t C_addr;
+} kernel_arg_t;
+
+#endif

tests/regression/sgemm_tpu/kernel.cpp

@@ -0,0 +1,24 @@
+#include <vx_spawn.h>
+#include "common.h"
+
+void kernel_body(kernel_arg_t* __UNIFORM__ arg) {
+	auto A = reinterpret_cast<TYPE*>(arg->A_addr);
+	auto B = reinterpret_cast<TYPE*>(arg->B_addr);
+	auto C = reinterpret_cast<TYPE*>(arg->C_addr);
+    auto size = arg->size;
+
+    int col = blockIdx.x;
+    int row = blockIdx.y;
+
+    TYPE sum(0);
+    for (int e = 0; e < size; ++e) {
+        sum += A[row * size + e] * B[e * size + col];
+    }
+
+    C[row * size + col] = sum;
+}
+
+int main() {
+	kernel_arg_t* arg = (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH);
+	return vx_spawn_threads(2, arg->grid_dim, nullptr, (vx_kernel_func_cb)kernel_body, arg);
+}

tests/regression/sgemm_tpu/main.cpp

@@ -0,0 +1,239 @@
+#include <iostream>
+#include <unistd.h>
+#include <string.h>
+#include <vector>
+#include <chrono>
+#include <vortex.h>
+#include <cmath>
+#include "common.h"
+
+#define FLOAT_ULP 6
+
+#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)
+
+///////////////////////////////////////////////////////////////////////////////
+
+template <typename Type>
+class Comparator {};
+
+template <>
+class Comparator<int> {
+public:
+  static const char* type_str() {
+    return "integer";
+  }
+  static int generate() {
+    return rand();
+  }
+  static bool compare(int a, int b, int index, int errors) {
+    if (a != b) {
+      if (errors < 100) {
+        printf("*** error: [%d] expected=%d, actual=%d\n", index, b, a);
+      }
+      return false;
+    }
+    return true;
+  }
+};
+
+template <>
+class Comparator<float> {
+public:
+  static const char* type_str() {
+    return "float";
+  }
+  static int generate() {
+    return static_cast<float>(rand()) / RAND_MAX;
+  }
+  static bool compare(float a, float b, int index, int errors) {
+    union fi_t { float f; int32_t i; };
+    fi_t fa, fb;
+    fa.f = a;
+    fb.f = b;
+    auto d = std::abs(fa.i - fb.i);
+    if (d > FLOAT_ULP) {
+      if (errors < 100) {
+        printf("*** error: [%d] expected=%f, actual=%f\n", index, b, a);
+      }
+      return false;
+    }
+    return true;
+  }
+};
+
+static void matmul_cpu(TYPE* out, const TYPE* A, const TYPE* B, uint32_t width, uint32_t height) {
+  for (uint32_t row = 0; row < height; ++row) {
+    for (uint32_t col = 0; col < width; ++col) {
+      TYPE sum(0);
+      for (uint32_t e = 0; e < width; ++e) {
+          sum += A[row * width + e] * B[e * width + col];
+      }
+      out[row * width + col] = sum;
+    }
+  }
+}
+
+const char* kernel_file = "kernel.vxbin";
+uint32_t size = 32;
+
+vx_device_h device = nullptr;
+vx_buffer_h A_buffer = nullptr;
+vx_buffer_h B_buffer = nullptr;
+vx_buffer_h C_buffer = nullptr;
+vx_buffer_h krnl_buffer = nullptr;
+vx_buffer_h args_buffer = nullptr;
+kernel_arg_t kernel_arg = {};
+
+static void show_usage() {
+   std::cout << "Vortex Test." << std::endl;
+   std::cout << "Usage: [-k: kernel] [-n size] [-h: help]" << std::endl;
+}
+
+static void parse_args(int argc, char **argv) {
+  int c;
+  while ((c = getopt(argc, argv, "n:k:h")) != -1) {
+    switch (c) {
+    case 'n':
+      size = atoi(optarg);
+      break;
+    case 'k':
+      kernel_file = optarg;
+      break;
+    case 'h':
+      show_usage();
+      exit(0);
+      break;
+    default:
+      show_usage();
+      exit(-1);
+    }
+  }
+}
+
+void cleanup() {
+  if (device) {
+    vx_mem_free(A_buffer);
+    vx_mem_free(B_buffer);
+    vx_mem_free(C_buffer);
+    vx_mem_free(krnl_buffer);
+    vx_mem_free(args_buffer);
+    vx_dev_close(device);
+  }
+}
+
+int main(int argc, char *argv[]) {
+  // parse command arguments
+  parse_args(argc, argv);
+
+  std::srand(50);
+
+  // open device connection
+  std::cout << "open device connection" << std::endl;
+  RT_CHECK(vx_dev_open(&device));
+
+  uint32_t size_sq = size * size;
+  uint32_t buf_size = size_sq * sizeof(TYPE);
+
+  std::cout << "data type: " << Comparator<TYPE>::type_str() << std::endl;
+  std::cout << "matrix size: " << size << "x" << size << std::endl;
+
+  kernel_arg.grid_dim[0] = size;
+  kernel_arg.grid_dim[1] = size;
+  kernel_arg.size = size;
+
+  // allocate device memory
+  std::cout << "allocate device memory" << std::endl;
+  RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_READ, &A_buffer));
+  RT_CHECK(vx_mem_address(A_buffer, &kernel_arg.A_addr));
+  RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_READ, &B_buffer));
+  RT_CHECK(vx_mem_address(B_buffer, &kernel_arg.B_addr));
+  RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_WRITE, &C_buffer));
+  RT_CHECK(vx_mem_address(C_buffer, &kernel_arg.C_addr));
+
+  std::cout << "A_addr=0x" << std::hex << kernel_arg.A_addr << std::endl;
+  std::cout << "B_addr=0x" << std::hex << kernel_arg.B_addr << std::endl;
+  std::cout << "C_addr=0x" << std::hex << kernel_arg.C_addr << std::endl;
+
+  // generate source data
+  std::vector<TYPE> h_A(size_sq);
+  std::vector<TYPE> h_B(size_sq);
+  std::vector<TYPE> h_C(size_sq);
+  for (uint32_t i = 0; i < size_sq; ++i) {
+    h_A[i] = Comparator<TYPE>::generate();
+    h_B[i] = Comparator<TYPE>::generate();
+  }
+
+  // upload matrix A buffer
+  {
+    std::cout << "upload matrix A buffer" << std::endl;
+    RT_CHECK(vx_copy_to_dev(A_buffer, h_A.data(), 0, buf_size));
+  }
+
+  // upload matrix B buffer
+  {
+    std::cout << "upload matrix B buffer" << std::endl;
+    RT_CHECK(vx_copy_to_dev(B_buffer, h_B.data(), 0, buf_size));
+  }
+
+  // upload program
+  std::cout << "upload program" << std::endl;
+  RT_CHECK(vx_upload_kernel_file(device, kernel_file, &krnl_buffer));
+
+  // upload kernel argument
+  std::cout << "upload kernel argument" << std::endl;
+  RT_CHECK(vx_upload_bytes(device, &kernel_arg, sizeof(kernel_arg_t), &args_buffer));
+
+  auto time_start = std::chrono::high_resolution_clock::now();
+
+  // start device
+  std::cout << "start device" << std::endl;
+  RT_CHECK(vx_start(device, krnl_buffer, args_buffer));
+
+  // wait for completion
+  std::cout << "wait for completion" << std::endl;
+  RT_CHECK(vx_ready_wait(device, VX_MAX_TIMEOUT));
+
+  auto time_end = std::chrono::high_resolution_clock::now();
+  double elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(time_end - time_start).count();
+  printf("Elapsed time: %lg ms\n", elapsed);
+
+  // download destination buffer
+  std::cout << "download destination buffer" << std::endl;
+  RT_CHECK(vx_copy_from_dev(h_C.data(), C_buffer, 0, buf_size));
+
+  // verify result
+  std::cout << "verify result" << std::endl;
+  int errors = 0;
+  {
+    std::vector<TYPE> h_ref(size_sq);
+    matmul_cpu(h_ref.data(), h_A.data(), h_B.data(), size, size);
+
+    for (uint32_t i = 0; i < h_ref.size(); ++i) {
+      if (!Comparator<TYPE>::compare(h_C[i], h_ref[i], i, errors)) {
+        ++errors;
+      }
+    }
+  }
+
+  // cleanup
+  std::cout << "cleanup" << std::endl;
+  cleanup();
+
+  if (errors != 0) {
+    std::cout << "Found " << std::dec << errors << " errors!" << std::endl;
+    std::cout << "FAILED!" << std::endl;
+    return errors;
+  }
+
+  std::cout << "PASSED!" << std::endl;
+
+  return 0;
+}