Bug/packld nanbox

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

tests/regression/packld_box/common.h

@@ -0,0 +1,18 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#include <stdint.h>
+
+// 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_

tests/regression/packld_box/kernel.cpp

@@ -0,0 +1,68 @@
+#include <vx_spawn2.h>
+#include <vx_intrinsics.h>
+#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<const uint8_t*>(arg->src_addr);
+    // uint64_t, not float: a float store would drop the NaN-box half
+    auto dst_lb   = reinterpret_cast<uint64_t*>(arg->dst_lb_addr);
+    auto dst_lh   = reinterpret_cast<uint64_t*>(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*/);
+    }
+}

tests/regression/packld_box/main.cpp

@@ -0,0 +1,192 @@
+#include <iostream>
+#include <unistd.h>
+#include <string.h>
+#include <cstdlib>
+#include <vector>
+#include <cstdint>
+#include <cmath>
+#include <vortex2.h>
+#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<uint8_t> 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<uint64_t> h_dst_lb(num_tasks * NUM_POINTS);
+  std::vector<uint64_t> 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;
+}