From 2d78a1e7c6df993337af886f74a88c4fd8480686 Mon Sep 17 00:00:00 2001 From: cassuto Date: Fri, 13 Feb 2026 07:12:50 -0500 Subject: [PATCH 1/2] add test for srai --- tests/regression/arith/Makefile | 14 +++ tests/regression/arith/common.h | 16 +++ tests/regression/arith/kernel.cpp | 14 +++ tests/regression/arith/main.cpp | 198 ++++++++++++++++++++++++++++++ 4 files changed, 242 insertions(+) create mode 100644 tests/regression/arith/Makefile create mode 100644 tests/regression/arith/common.h create mode 100644 tests/regression/arith/kernel.cpp create mode 100644 tests/regression/arith/main.cpp diff --git a/tests/regression/arith/Makefile b/tests/regression/arith/Makefile new file mode 100644 index 0000000000..b671c2837a --- /dev/null +++ b/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 \ No newline at end of file diff --git a/tests/regression/arith/common.h b/tests/regression/arith/common.h new file mode 100644 index 0000000000..dc68f613af --- /dev/null +++ b/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 diff --git a/tests/regression/arith/kernel.cpp b/tests/regression/arith/kernel.cpp new file mode 100644 index 0000000000..7a07bfb34e --- /dev/null +++ b/tests/regression/arith/kernel.cpp @@ -0,0 +1,14 @@ +#include +#include "common.h" + +void kernel_body(kernel_arg_t* __UNIFORM__ arg) { + auto src0_ptr = reinterpret_cast(arg->src0_addr); + auto dst_ptr = reinterpret_cast(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); +} diff --git a/tests/regression/arith/main.cpp b/tests/regression/arith/main.cpp new file mode 100644 index 0000000000..093cee77d1 --- /dev/null +++ b/tests/regression/arith/main.cpp @@ -0,0 +1,198 @@ +#include +#include +#include +#include +#include +#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 +class Comparator {}; + + +template <> +class Comparator { +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 { +public: + static const char* type_str() { + return "integer 64bit"; + } + static int64_t generate() { + return (static_cast(rand()) << 32) | static_cast(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_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 h_src0(num_points); + std::vector h_dst(num_points); + + for (uint32_t i = 0; i < num_points; ++i) { + h_src0[i] = (1ul<<63) | Comparator::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::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; +} \ No newline at end of file From 7fd04799d16744f4e3919bb7b38df489a9b5acd2 Mon Sep 17 00:00:00 2001 From: cassuto Date: Fri, 13 Feb 2026 07:13:01 -0500 Subject: [PATCH 2/2] fix srai decode --- sim/simx/decode.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sim/simx/decode.cpp b/sim/simx/decode.cpp index f2d6cc355c..4a7ba647d1 100644 --- a/sim/simx/decode.cpp +++ b/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