Fix srai

sim/simx/decode.cpp

@@ -624,7 +624,7 @@ void Emulator::decode(uint32_t code, uint32_t wid, uint64_t uuid) {
         break;
       }
       case 5: { // RV32I: SRA/SRL
-        instr->setOpType((funct7 == 0x20) ? AluType::SRA : AluType::SRL);
+        instr->setOpType((funct7 & 0x20) ? AluType::SRA : AluType::SRL);
         break;
       }
       case 6: { // RV32I: OR

tests/regression/arith/Makefile

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

tests/regression/arith/common.h

@@ -0,0 +1,16 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#ifndef TYPE
+#define TYPE int64_t
+#endif
+
+#define SHIFT_BITS 47
+
+typedef struct {
+  uint32_t num_points;
+  uint64_t src0_addr;
+  uint64_t dst_addr;  
+} kernel_arg_t;
+
+#endif

tests/regression/arith/kernel.cpp

@@ -0,0 +1,14 @@
+#include <vx_spawn.h>
+#include "common.h"
+
+void kernel_body(kernel_arg_t* __UNIFORM__ arg) {
+	auto src0_ptr = reinterpret_cast<TYPE*>(arg->src0_addr);
+	auto dst_ptr  = reinterpret_cast<TYPE*>(arg->dst_addr);
+
+	dst_ptr[blockIdx.x] = src0_ptr[blockIdx.x] >> SHIFT_BITS;
+}
+
+int main() {
+	kernel_arg_t* arg = (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH);
+	return vx_spawn_threads(1, &arg->num_points, nullptr, (vx_kernel_func_cb)kernel_body, arg);
+}

tests/regression/arith/main.cpp

@@ -0,0 +1,198 @@
+#include <iostream>
+#include <unistd.h>
+#include <string.h>
+#include <vector>
+#include <vortex.h>
+#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<int32_t> {
+public:
+  static const char* type_str() {
+    return "integer 32bit";
+  }
+  static int32_t generate() {
+    return  rand();
+  }
+  static bool compare(int32_t a, int32_t b, int index, int errors) {
+    if (a != b) {
+      if (errors < 100) {
+        printf("*** error: [%d] expected=%x, actual=%x\n", index, b, a);
+      }
+      return false;
+    }
+    return true;
+  }
+};
+
+template <>
+class Comparator<int64_t> {
+public:
+  static const char* type_str() {
+    return "integer 64bit";
+  }
+  static int64_t generate() {
+    return (static_cast<int64_t>(rand()) << 32) | static_cast<int64_t>(rand());
+  }
+  static bool compare(int64_t a, int64_t b, int index, int errors) {
+    if (a != b) {
+      if (errors < 100) {
+        printf("*** error: [%d] expected=%lx, actual=%lx\n", index, b, a);
+      }
+      return false;
+    }
+    return true;
+  }
+};
+
+const char* kernel_file = "kernel.vxbin";
+uint32_t size = 16;
+
+vx_device_h device = nullptr;
+vx_buffer_h src0_buffer = nullptr;
+vx_buffer_h dst_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 words] [-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(src0_buffer);
+    vx_mem_free(dst_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 num_points = size;
+  uint32_t buf_size = num_points * sizeof(TYPE);
+
+  std::cout << "number of points: " << num_points << std::endl;
+  std::cout << "data type: " << Comparator<TYPE>::type_str() << std::endl;
+  std::cout << "buffer size: " << buf_size << " bytes" << std::endl;
+
+  kernel_arg.num_points = num_points;
+
+  // allocate device memory
+  std::cout << "allocate device memory" << std::endl;
+  RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_READ, &src0_buffer));
+  RT_CHECK(vx_mem_address(src0_buffer, &kernel_arg.src0_addr));
+  RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_WRITE, &dst_buffer));
+  RT_CHECK(vx_mem_address(dst_buffer, &kernel_arg.dst_addr));
+
+  std::cout << "dev_src0=0x" << std::hex << kernel_arg.src0_addr << std::endl;
+  std::cout << "dev_dst=0x" << std::hex << kernel_arg.dst_addr << std::endl;
+
+  // allocate host buffers
+  std::cout << "allocate host buffers" << std::endl;
+  std::vector<TYPE> h_src0(num_points);
+  std::vector<TYPE> h_dst(num_points);
+
+  for (uint32_t i = 0; i < num_points; ++i) {
+    h_src0[i] = (1ul<<63) | Comparator<TYPE>::generate();
+  }
+
+  // upload source buffer0
+  std::cout << "upload source buffer0" << std::endl;
+  RT_CHECK(vx_copy_to_dev(src0_buffer, h_src0.data(), 0, buf_size));
+
+  // Upload kernel binary
+  std::cout << "Upload kernel binary" << 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));
+
+  // 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));
+
+  // download destination buffer
+  std::cout << "download destination buffer" << std::endl;
+  RT_CHECK(vx_copy_from_dev(h_dst.data(), dst_buffer, 0, buf_size));
+
+  // verify result
+  std::cout << "verify result" << std::endl;
+  int errors = 0;
+  for (uint32_t i = 0; i < num_points; ++i) {
+    auto ref = h_src0[i] >> SHIFT_BITS;
+    auto cur = h_dst[i];
+    if (!Comparator<TYPE>::compare(cur, ref, 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 1;
+  }
+
+  std::cout << "PASSED!" << std::endl;
+
+  return 0;
+}