vecadd working:

fix bugs from merging kmu onto bug fixes

Author: Randy L

hw/VX_config.toml

@@ -32,7 +32,7 @@ VM_ENABLED = "expr: 1 if $VM_ENABLE else 0"
 EXT_D_ENABLE = "expr: $XLEN_64"
 FLEN = "expr: 64 if $EXT_D_ENABLE else 32"
 
-KMU_ENABLE = false
+KMU_ENABLE = true
 KMU_ENABLED = "expr: 1 if $KMU_ENABLE else 0"
 
 # extensions

kernel/include/vx_spawn.h

@@ -15,6 +15,7 @@
 #define __VX_SPAWN_H__
 
 #include <vx_intrinsics.h>
+#include <vx_print.h>
 #include <stdint.h>
 
 #ifdef __cplusplus
@@ -30,10 +31,7 @@ typedef union {
   uint32_t m[3];
 } dim3_t;
 
-extern __thread dim3_t blockIdx;
-extern __thread dim3_t threadIdx;
 extern dim3_t gridDim;
-extern dim3_t blockDim;
 
 extern __thread uint32_t __local_group_id;
 extern uint32_t __warps_per_group;

kernel/src/vx_start.S

@@ -23,22 +23,8 @@ _start:
 
 #ifdef KMU_ENABLE
 
-  # set global pointer register
-  la    gp, __global_pointer
-
-  # set stack pointer register
-  LOAD_IMMEDIATE64(sp, STACK_BASE_ADDR)
-  csrr  t0, VX_CSR_MHARTID
-  sll   t1, t0, STACK_LOG2_SIZE
-  sub   sp, sp, t1
-
-  # set thread pointer register
-  # use address space after BSS region
-  la    t1, __tbss_size
-  mul   t0, t0, t1
-  la    tp, _end
-  add   tp, tp, t0
-  ret
+  # initialize per-thread registers
+  jal   init_regs
 
   # initialize TLS for all warps
   call  __init_tls
@@ -53,12 +39,8 @@ _start:
 #endif
 
   # call main program routine
-  csrr  a0, VX_CSR_MSCRATCH
   call  main
 
-  # Moving to single thread
-  # vx_tmc(0 == vx_warp_id());
-
   # call exit routine
   tail  exit
 .size _start, .-_start
@@ -108,6 +90,8 @@ _start:
   tail  exit
 .size _start, .-_start
 
+#endif
+
 .section .text
 .type init_regs, @function
 .local init_regs
@@ -132,6 +116,8 @@ init_regs:
   add   tp, tp, t0
   ret
 
+#ifndef KMU_ENABLE
+
 .section .text
 .type init_regs_all, @function
 .local init_regs_all

riscv-openocd

@@ -0,0 +1 @@
+Subproject commit eb01c632a4bb1c07d2bddb008d6987c809f1c496

runtime/rtlsim/vortex.cpp

@@ -20,6 +20,7 @@
 #include <stdint.h>
 #include <stdio.h>
 #include <stdlib.h>
+#include <string.h>
 #include <assert.h>
 #include <iostream>
 #include <future>
@@ -211,6 +212,27 @@ class vx_device {
     this->dcr_write(VX_DCR_BASE_STARTUP_ARG0, args_addr & 0xffffffff);
     this->dcr_write(VX_DCR_BASE_STARTUP_ARG1, args_addr >> 32);
 
+    // read block and grid dimensions from kernel arguments
+    uint32_t block_dim[3] = {1, 1, 1};
+    uint32_t grid_dim[3] = {1, 1, 1};
+    if (args_addr != 0) {
+      // Read first 24 bytes of kernel arguments (block_dim[3] + grid_dim[3])
+      struct {
+        uint32_t block_dim[3];
+        uint32_t grid_dim[3];
+      } args_dims;
+      this->download(&args_dims, args_addr, sizeof(args_dims));
+      memcpy(block_dim, args_dims.block_dim, sizeof(block_dim));
+      memcpy(grid_dim, args_dims.grid_dim, sizeof(grid_dim));
+    }
+
+    this->dcr_write(VX_DCR_BASE_GRID_DIM0, grid_dim[0]);
+    this->dcr_write(VX_DCR_BASE_GRID_DIM1, grid_dim[1]);
+    this->dcr_write(VX_DCR_BASE_GRID_DIM2, grid_dim[2]);
+    this->dcr_write(VX_DCR_BASE_BLOCK_DIM0, block_dim[0]);
+    this->dcr_write(VX_DCR_BASE_BLOCK_DIM1, block_dim[1]);
+    this->dcr_write(VX_DCR_BASE_BLOCK_DIM2, block_dim[2]);
+
     // start new run
     future_ = std::async(std::launch::async, [&]{
       processor_.run();

tests/regression/vecadd/common.h

@@ -6,6 +6,8 @@
 #endif
 
 typedef struct {
+  uint32_t block_dim[3];
+  uint32_t grid_dim[3];
   uint32_t num_points;
   uint64_t src0_addr;
   uint64_t src1_addr;

tests/regression/vecadd/kernel.cpp

@@ -42,10 +42,13 @@ int main() {
 	// Calculate global thread ID
 	// threadIdx.x gives the flat thread index within the CTA (warp_local_id * NUM_THREADS + thread_id)
 	// globalId = blockIdx.x * blockDim.x + threadIdx.x
-	uint32_t globalId = blockIdx.x * blockDim.x + threadIdx.x;
+	uint32_t bidx = blockIdx.x;
+	uint32_t tidx = threadIdx.x;
+	uint32_t bdim = blockDim.x;
+	uint32_t globalId = bidx * bdim + tidx;
 
 	vx_printf("block id x: %d, threadIdx.x: %d, global id: %d\n", 
-		blockIdx.x, threadIdx.x, globalId);
+		bidx, tidx, globalId);
 
 	auto src0_ptr = reinterpret_cast<TYPE*>(arg->src0_addr);
 	auto src1_ptr = reinterpret_cast<TYPE*>(arg->src1_addr);

tests/regression/vecadd/main.cpp

@@ -134,6 +134,17 @@ int main(int argc, char *argv[]) {
   std::cout << "data type: " << Comparator<TYPE>::type_str() << std::endl;
   std::cout << "buffer size: " << buf_size << " bytes" << std::endl;
 
+  // Set block dimensions to 8 (warp width)
+  const uint32_t block_size = 16;
+  kernel_arg.block_dim[0] = block_size;
+  kernel_arg.block_dim[1] = 1;
+  kernel_arg.block_dim[2] = 1;
+
+  // Calculate grid dimensions based on number of points
+  kernel_arg.grid_dim[0] = (num_points + block_size - 1) / block_size;  // Round up
+  kernel_arg.grid_dim[1] = 1;
+  kernel_arg.grid_dim[2] = 1;
+
   kernel_arg.num_points = num_points;
 
   // allocate device memory