Added test for vote

tests/regression/vote/Makefile

@@ -0,0 +1,14 @@
+ROOT_DIR := $(realpath ../../..)
+include $(ROOT_DIR)/config.mk
+
+PROJECT := vote
+
+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/vote/common.h

@@ -0,0 +1,12 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+typedef struct {
+  uint32_t num_tasks;
+  uint32_t task_size;
+  uint64_t src0_addr;
+  uint64_t src1_addr;
+  uint64_t dst_addr;
+} kernel_arg_t;
+
+#endif

tests/regression/vote/kernel.cpp

@@ -0,0 +1,35 @@
+#include <vx_spawn.h>
+#include "common.h"
+#include <vx_intrinsics.h>
+
+void kernel_body(kernel_arg_t* __UNIFORM__ arg) {
+	uint32_t count    = arg->task_size;
+	int32_t* src0_ptr = (int32_t*)arg->src0_addr;
+	int32_t* src1_ptr = (int32_t*)arg->src1_addr;
+	int32_t* dst_ptr  = (int32_t*)arg->dst_addr;
+
+	uint32_t offset = blockIdx.x * count;
+	for (uint32_t i = 0; i < count; ++i) {
+		dst_ptr[offset+i] = src0_ptr[offset+i] + src1_ptr[offset+i];
+	} 
+	vx_tmc(1);
+	int val = 14;
+	vx_store(val,3);
+	vx_tmc(2);
+	val = 5;
+	vx_store(val,3); 
+	vx_tmc(4);
+	val =45;     
+	vx_store(val,3); 
+	vx_tmc(8);
+	val = 120;
+	vx_store(val,3);         
+	vx_tmc(15); 
+	vx_vote();       
+}                
+          
+int main() {
+	kernel_arg_t* arg = (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH);
+	return vx_spawn_threads(1, &arg->num_tasks, nullptr, (vx_kernel_func_cb)kernel_body, arg);
+}
+      

tests/regression/vote/main.cpp

@@ -0,0 +1,175 @@
+#include <iostream>
+#include <unistd.h>
+#include <string.h>
+#include <vector>
+#include <vortex.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";
+uint32_t count = 0;
+
+vx_device_h device = nullptr;
+vx_buffer_h src0_buffer = nullptr;
+vx_buffer_h src1_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':
+      count = atoi(optarg);
+      break;
+    case 'k':
+      kernel_file = optarg;
+      break;
+    case 'h':
+    case '?': {
+      show_usage();
+      exit(0);
+    } break;
+    default:
+      show_usage();
+      exit(-1);
+    }
+  }
+}
+
+void cleanup() {
+  if (device) {
+    vx_mem_free(src0_buffer);
+    vx_mem_free(src1_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);
+
+  if (count == 0) {
+    count = 1;
+  }
+
+  // open device connection
+  std::cout << "open device connection" << std::endl;
+  RT_CHECK(vx_dev_open(&device));
+
+  uint64_t num_cores, num_warps, num_threads;
+  RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_CORES, &num_cores));
+  RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_WARPS, &num_warps));
+  RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_THREADS, &num_threads));
+
+  uint32_t total_threads = num_cores * num_warps * num_threads;
+  uint32_t num_points = count * total_threads;
+  uint32_t buf_size   = num_points * sizeof(int32_t);
+
+  std::cout << "number of points: " << num_points << std::endl;
+  std::cout << "buffer size: " << buf_size << " bytes" << std::endl;
+
+  kernel_arg.num_tasks = total_threads;
+  kernel_arg.task_size = count;
+
+  // 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_READ, &src1_buffer));
+  RT_CHECK(vx_mem_address(src1_buffer, &kernel_arg.src1_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_src1=0x" << std::hex << kernel_arg.src1_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<int32_t> h_src0(num_points);
+  std::vector<int32_t> h_src1(num_points);
+  std::vector<int32_t> h_dst(num_points);
+
+  // generate source data
+  for (uint32_t i = 0; i < num_points; ++i) {
+    h_src0[i] = i-1;
+    h_src1[i] = i+1;
+  }
+
+  // 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 source buffer1
+  std::cout << "upload source buffer1" << std::endl;
+  RT_CHECK(vx_copy_to_dev(src1_buffer, h_src1.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));
+
+  // 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) {
+    int ref = i + i;
+    int cur = h_dst[i];
+    if (cur != ref) {
+      std::cout << "error at result #" << std::dec << i
+                << std::hex << ": actual 0x" << cur << ", expected 0x" << ref << std::endl;
+      ++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;
+}