From 23e85956a8dd31e8381a0a6c1451e1e3446dd107 Mon Sep 17 00:00:00 2001 From: RunjiaChen Date: Wed, 12 Aug 2026 15:02:36 +0800 Subject: [PATCH 1/2] =?UTF-8?q?tests/regression:=20packld=5Fbox=20?= =?UTF-8?q?=E2=80=94=20expose=20missing=20NaN-boxing=20on=20packed=20loads?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit vx_packlb_f / vx_packlh_f write a 32-bit float into a 64-bit FLEN register, so RISC-V requires the upper word to be NaN-boxed (all ones). SimX does this; the RTL does not, and the existing packld test cannot see the difference. The blindness is structural, not a weak assertion. kernel.cpp stores each result through a float*, which is a 32-bit fsw, and main.cpp compares four bytes. The NaN-box lives in the upper half and is discarded before it ever reaches memory, so both drivers pass. packld_box is a copy of packld with the observation window widened to the full 64-bit container: - kernel.cpp stores uint64_t instead of float, and reads the packed result with fmv.x.d fused into the same asm block as the packed load (as separate statements the compiler may spill the float through a 32-bit fsw/flw pair, which would fabricate the very bits under test); - main.cpp sizes the destination buffers for uint64_t and counts the packed VALUE and the NaN-BOX as separate error classes, so a box-only failure cannot be confused with a mis-packed value. Measured, upstream tooling only, default config (1 cluster, 1 core, 4 warps, 4 threads; EXT_D on so FLEN=64 and boxing is required): ci/blackbox.sh --driver=simx --app=packld_box -> PASSED value errors = 0, nanbox errors = 0 ci/blackbox.sh --driver=rtlsim --app=packld_box -> FAILED value errors = 0, nanbox errors = 512 value errors = 0 on both drivers: the packing itself (byte order, strided addressing, both widths) is correct everywhere. The NaN-box is the sole divergence, and it fails on every one of the 512 packed loads. The upper half is not merely un-set to ones, it is never written at all, so it retains prior register content: one observed sample read 0x8d2428b1 rather than zero, which makes the architectural result of a packed load depend on whatever the destination register happened to hold. Root cause is in the RTL, not the test. VX_decode.sv gives packed loads op_type INST_LSU_LBU / INST_LSU_LHU, VX_uop_packld.sv preserves that op_type while rewriting only offset and bytesel, and VX_lsu_slice.sv applies NaN-boxing solely on the LSU_FMT_W arm. The BU/HU arms are plain zero-extends, so rsp_is_float is computed but never consulted for these instructions. VX_uop_packld.sv's comment claiming the boxing "is handled automatically ... triggers nan-box in lsu_slice" does not hold. This change is purely additive: no existing file is modified. Adding packld_box to the tests/regression/Makefile list is deliberately omitted so nothing upstream changes; blackbox.sh --app=packld_box works without it. --- tests/regression/packld_box/Makefile | 16 +++ tests/regression/packld_box/common.h | 18 +++ tests/regression/packld_box/kernel.cpp | 60 ++++++++ tests/regression/packld_box/main.cpp | 192 +++++++++++++++++++++++++ 4 files changed, 286 insertions(+) create mode 100644 tests/regression/packld_box/Makefile create mode 100644 tests/regression/packld_box/common.h create mode 100644 tests/regression/packld_box/kernel.cpp create mode 100644 tests/regression/packld_box/main.cpp diff --git a/tests/regression/packld_box/Makefile b/tests/regression/packld_box/Makefile new file mode 100644 index 0000000000..3cb2003390 --- /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 0000000000..6f7f7c4c34 --- /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 0000000000..4cbc0ff3e4 --- /dev/null +++ b/tests/regression/packld_box/kernel.cpp @@ -0,0 +1,60 @@ +#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. +__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" + "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" + "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 0000000000..944fdbfa5b --- /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; +} From df41af4e6c35b6b19e6dcf79fb121e6ae4f2369c Mon Sep 17 00:00:00 2001 From: RunjiaChen Date: Wed, 12 Aug 2026 22:06:38 +0800 Subject: [PATCH 2/2] added some explanations for how the bug stays hidden --- tests/regression/packld_box/kernel.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/tests/regression/packld_box/kernel.cpp b/tests/regression/packld_box/kernel.cpp index 4cbc0ff3e4..b407b12518 100644 --- a/tests/regression/packld_box/kernel.cpp +++ b/tests/regression/packld_box/kernel.cpp @@ -10,11 +10,17 @@ // 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" ); @@ -26,6 +32,8 @@ 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" );