diff --git a/tests/regression/packld_box/Makefile b/tests/regression/packld_box/Makefile new file mode 100644 index 000000000..3cb200339 --- /dev/null +++ b/tests/regression/packld_box/Makefile @@ -0,0 +1,16 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +PROJECT := packld_box + +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 diff --git a/tests/regression/packld_box/common.h b/tests/regression/packld_box/common.h new file mode 100644 index 000000000..6f7f7c4c3 --- /dev/null +++ b/tests/regression/packld_box/common.h @@ -0,0 +1,18 @@ +#ifndef _COMMON_H_ +#define _COMMON_H_ + +#include + +// Number of test vectors per thread +#ifndef NUM_POINTS +#define NUM_POINTS 16 +#endif + +typedef struct { + uint32_t num_tasks; // total thread count + uint64_t src_addr; // byte array: 4*num_tasks*NUM_POINTS elements + uint64_t dst_lb_addr; // float output for vx_packlb_f: num_tasks*NUM_POINTS floats + uint64_t dst_lh_addr; // float output for vx_packlh_f: num_tasks*NUM_POINTS floats +} kernel_arg_t; + +#endif // _COMMON_H_ diff --git a/tests/regression/packld_box/kernel.cpp b/tests/regression/packld_box/kernel.cpp new file mode 100644 index 000000000..b407b1251 --- /dev/null +++ b/tests/regression/packld_box/kernel.cpp @@ -0,0 +1,68 @@ +#include +#include +#include "common.h" + +// Read the FULL 64-bit FP container produced by a packed load. +// +// The packed load and the fmv.x.d are fused into ONE asm block using an +// explicit FP register. As two separate statements the compiler is free to +// spill the float through a 32-bit fsw/flw pair, which would destroy (or +// silently recreate) the upper half -- exactly the bits under test. The +// instruction encoding is identical to vx_intrinsics.h's vx_packlb_f / +// vx_packlh_f: custom0, funct7=4, funct3=1/2. + +//The inline functions are needed to prevent register spilling. +// Register spilling silently erases the bug by doing proper NaN-boxing. +// Uncomment the two lines to see all the tests pass. +__attribute__((always_inline)) +inline uint64_t packlb_bits(const void* base, uint32_t stride) { + uint64_t bits; + __asm__ volatile ( + ".insn r %1, 1, 4, ft0, %2, %3\n\t" + //"fsw ft0, 12(sp)\n\t" / + //"flw ft0, 12(sp)\n\t" + "fmv.x.d %0, ft0" + : "=r"(bits) : "i"(RISCV_CUSTOM0), "r"(base), "r"(stride) : "ft0", "memory" + ); + return bits; +} + +__attribute__((always_inline)) +inline uint64_t packlh_bits(const void* base, uint32_t stride) { + uint64_t bits; + __asm__ volatile ( + ".insn r %1, 2, 4, ft0, %2, %3\n\t" + //"fsw ft0, 12(sp)\n\t" + //"flw ft0, 12(sp)\n\t" + "fmv.x.d %0, ft0" + : "=r"(bits) : "i"(RISCV_CUSTOM0), "r"(base), "r"(stride) : "ft0", "memory" + ); + return bits; +} + +// Each thread exercises vx_packlb_f and vx_packlh_f over NUM_POINTS vectors. +// Layout of src (byte array): +// For point p in thread t: src[t*4*NUM_POINTS + p*4 + lane] (stride = 1 byte) +// For PACKLB: base = &src[t*4*NUM_POINTS + p*4], stride = 1 +// → result = b0 | (b1<<8) | (b2<<16) | (b3<<24) +// For PACKLH: base = &src_u16[t*2*NUM_POINTS + p*2], stride = 2 bytes +// → result = h0 | (h1<<16) +__kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) { + auto src_ptr = reinterpret_cast(arg->src_addr); + // uint64_t, not float: a float store would drop the NaN-box half + auto dst_lb = reinterpret_cast(arg->dst_lb_addr); + auto dst_lh = reinterpret_cast(arg->dst_lh_addr); + + uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t stride = 1; // byte stride between consecutive elements + + for (uint32_t p = 0; p < NUM_POINTS; ++p) { + // 4 bytes at consecutive addresses → one packed float (PACKLB) + const uint8_t* base_lb = src_ptr + (tid * 4 * NUM_POINTS + p * 4); + dst_lb[tid * NUM_POINTS + p] = packlb_bits(base_lb, stride); + + // 2 halfwords at consecutive addresses → one packed float (PACKLH) + const uint8_t* base_lh = src_ptr + (tid * 4 * NUM_POINTS + p * 4); + dst_lh[tid * NUM_POINTS + p] = packlh_bits(base_lh, 2 /*halfword stride*/); + } +} diff --git a/tests/regression/packld_box/main.cpp b/tests/regression/packld_box/main.cpp new file mode 100644 index 000000000..944fdbfa5 --- /dev/null +++ b/tests/regression/packld_box/main.cpp @@ -0,0 +1,192 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include "common.h" + +#define RT_CHECK(_expr) \ + do { \ + int _ret = (_expr); \ + if (0 == _ret) break; \ + printf("Error: '%s' returned %d!\n", #_expr, _ret); \ + cleanup(); \ + exit(-1); \ + } while (false) + +static const char* kernel_file = "kernel.vxbin"; + +static vx_device_h device = nullptr; +static vx_buffer_h src_buffer = nullptr; +static vx_buffer_h dst_lb_buf = nullptr; +static vx_buffer_h dst_lh_buf = nullptr; +static vx_queue_h queue = nullptr; +static vx_module_h module_ = nullptr; +static vx_kernel_h kernel = nullptr; +static kernel_arg_t kernel_arg = {}; + +static void show_usage() { + std::cout << "Vortex packld test.\n" + << "Usage: [-k kernel] [-h help]\n"; +} + +static void parse_args(int argc, char** argv) { + int c; + while ((c = getopt(argc, argv, "k:h")) != -1) { + switch (c) { + case 'k': kernel_file = optarg; break; + case 'h': show_usage(); exit(0); + default: show_usage(); exit(-1); + } + } +} + +static void cleanup() { + if (device) { + if (src_buffer) vx_buffer_release(src_buffer); + if (dst_lb_buf) vx_buffer_release(dst_lb_buf); + if (dst_lh_buf) vx_buffer_release(dst_lh_buf); + if (kernel) vx_kernel_release(kernel); + if (module_) vx_module_release(module_); + if (queue) vx_queue_release(queue); + vx_device_dump_perf(device, stdout); + vx_device_release(device); + } +} + +int main(int argc, char* argv[]) { + parse_args(argc, argv); + std::srand(42); + + std::cout << "open device connection\n"; + RT_CHECK(vx_device_open(0, &device)); + + vx_queue_info_t qi = { sizeof(qi), nullptr, VX_QUEUE_PRIORITY_NORMAL, 0 }; + RT_CHECK(vx_queue_create(device, &qi, &queue)); + + uint64_t num_cores, num_warps, num_threads; + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_CORES, &num_cores)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_WARPS, &num_warps)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_THREADS, &num_threads)); + + uint32_t num_tasks = (uint32_t)(num_cores * num_warps * num_threads); + + // src layout: num_tasks * NUM_POINTS * 4 bytes + uint32_t src_bytes = num_tasks * NUM_POINTS * 4; + uint32_t dst_bytes = num_tasks * NUM_POINTS * sizeof(uint64_t); + + std::cout << "num_tasks=" << num_tasks + << " src=" << src_bytes << "B" + << " dst_lb=" << dst_bytes << "B" + << " dst_lh=" << dst_bytes << "B\n"; + + kernel_arg.num_tasks = num_tasks; + + // allocate device memory + RT_CHECK(vx_buffer_create(device, src_bytes, VX_MEM_READ, &src_buffer)); + RT_CHECK(vx_buffer_create(device, dst_bytes, VX_MEM_WRITE, &dst_lb_buf)); + RT_CHECK(vx_buffer_create(device, dst_bytes, VX_MEM_WRITE, &dst_lh_buf)); + RT_CHECK(vx_buffer_address(src_buffer, &kernel_arg.src_addr)); + RT_CHECK(vx_buffer_address(dst_lb_buf, &kernel_arg.dst_lb_addr)); + RT_CHECK(vx_buffer_address(dst_lh_buf, &kernel_arg.dst_lh_addr)); + + // generate random byte source data + std::vector h_src(src_bytes); + for (auto& b : h_src) b = (uint8_t)(std::rand() & 0xFF); + + RT_CHECK(vx_enqueue_write(queue, src_buffer, 0, h_src.data(), src_bytes, 0, nullptr, nullptr)); + + // load kernel module + RT_CHECK(vx_module_load_file(device, kernel_file, &module_)); + RT_CHECK(vx_module_get_kernel(module_, "main", &kernel)); + + std::cout << "start device\n"; + vx_event_h launch_ev = nullptr, read_ev_lb = nullptr, read_ev_lh = nullptr; + uint32_t grid_dim[1], block_dim[1]; + RT_CHECK(vx_device_max_occupancy_grid(device, 1, &num_tasks, grid_dim, block_dim)); + { + vx_launch_info_t li = {}; + li.struct_size = sizeof(li); + li.kernel = kernel; + li.args_host = &kernel_arg; + li.args_size = sizeof(kernel_arg); + li.ndim = 1; + li.grid_dim[0] = grid_dim[0]; + li.block_dim[0] = block_dim[0]; + RT_CHECK(vx_enqueue_launch(queue, &li, 0, nullptr, &launch_ev)); + } + + // download results + std::vector h_dst_lb(num_tasks * NUM_POINTS); + std::vector h_dst_lh(num_tasks * NUM_POINTS); + RT_CHECK(vx_enqueue_read(queue, h_dst_lb.data(), dst_lb_buf, 0, dst_bytes, 1, &launch_ev, &read_ev_lb)); + RT_CHECK(vx_enqueue_read(queue, h_dst_lh.data(), dst_lh_buf, 0, dst_bytes, 1, &launch_ev, &read_ev_lh)); + RT_CHECK(vx_event_wait_value(read_ev_lh, 1, VX_TIMEOUT_INFINITE)); + vx_event_release(read_ev_lh); + vx_event_release(read_ev_lb); + vx_event_release(launch_ev); + + // verify. The packed VALUE (low 32 bits) and the NaN-BOX (upper 32 bits) + // are counted as separate error classes, so a box-only failure is + // unambiguous and cannot be confused with a mis-packed value. + int value_errors = 0; + int box_errors = 0; + for (uint32_t t = 0; t < num_tasks; ++t) { + for (uint32_t p = 0; p < NUM_POINTS; ++p) { + uint32_t base_off = (t * NUM_POINTS + p) * 4; + uint32_t idx = t * NUM_POINTS + p; + + // PACKLB: pack 4 bytes with stride 1 + uint32_t ref_lb = (uint32_t)h_src[base_off + 0] + | ((uint32_t)h_src[base_off + 1] << 8) + | ((uint32_t)h_src[base_off + 2] << 16) + | ((uint32_t)h_src[base_off + 3] << 24); + uint64_t got_lb = h_dst_lb[idx]; + if ((uint32_t)got_lb != ref_lb) { + if (value_errors < 4) + printf("PACKLB VALUE error t=%u p=%u: expected=0x%08x got=0x%08x\n", + t, p, ref_lb, (uint32_t)got_lb); + ++value_errors; + } + if ((uint32_t)(got_lb >> 32) != 0xffffffffu) { + if (box_errors < 4) + printf("PACKLB NANBOX error t=%u p=%u: upper32 expected=0xffffffff got=0x%08x (full=0x%016llx)\n", + t, p, (uint32_t)(got_lb >> 32), (unsigned long long)got_lb); + ++box_errors; + } + + // PACKLH: pack 2 halfwords with stride 2 + uint16_t h0, h1; + memcpy(&h0, &h_src[base_off + 0], 2); + memcpy(&h1, &h_src[base_off + 2], 2); + uint32_t ref_lh = (uint32_t)h0 | ((uint32_t)h1 << 16); + uint64_t got_lh = h_dst_lh[idx]; + if ((uint32_t)got_lh != ref_lh) { + if (value_errors < 4) + printf("PACKLH VALUE error t=%u p=%u: expected=0x%08x got=0x%08x\n", + t, p, ref_lh, (uint32_t)got_lh); + ++value_errors; + } + if ((uint32_t)(got_lh >> 32) != 0xffffffffu) { + if (box_errors < 4) + printf("PACKLH NANBOX error t=%u p=%u: upper32 expected=0xffffffff got=0x%08x (full=0x%016llx)\n", + t, p, (uint32_t)(got_lh >> 32), (unsigned long long)got_lh); + ++box_errors; + } + } + } + + cleanup(); + + printf("value errors = %d, nanbox errors = %d\n", value_errors, box_errors); + int errors = value_errors + box_errors; + if (errors != 0) { + std::cout << "Found " << errors << " errors!\nFAILED!\n"; + return 1; + } + std::cout << "PASSED!\n"; + return 0; +}