Merge pull request #155 from JacobLevinson/stencil3d

Added Stencil3d regression test

.gitignore

@@ -1,2 +1,3 @@
 /build*
-/.vscode
+/.vscode
+*.cache

tests/regression/Makefile

@@ -15,6 +15,7 @@ all:
 	$(MAKE) -C sgemmx
 	$(MAKE) -C conv3x
 	$(MAKE) -C sgemm2x
+	$(MAKE) -C stencil3d
 
 run-simx:
 	$(MAKE) -C basic run-simx
@@ -30,6 +31,7 @@ run-simx:
 	$(MAKE) -C sgemmx run-simx
 	$(MAKE) -C conv3x run-simx
 	$(MAKE) -C sgemm2x run-simx
+	$(MAKE) -C stencil3d run-simx
 
 run-rtlsim:
 	$(MAKE) -C basic run-rtlsim
@@ -45,6 +47,7 @@ run-rtlsim:
 	$(MAKE) -C sgemmx run-rtlsim
 	$(MAKE) -C conv3x run-rtlsim
 	$(MAKE) -C sgemm2x run-rtlsim
+	$(MAKE) -C stencil3d run-rtlsim
 
 clean:
 	$(MAKE) -C basic clean
@@ -59,4 +62,5 @@ clean:
 	$(MAKE) -C vecaddx clean
 	$(MAKE) -C sgemmx clean
 	$(MAKE) -C conv3x clean
-	$(MAKE) -C sgemm2x clean
+	$(MAKE) -C sgemm2x clean
+	$(MAKE) -C stencil3d clean

tests/regression/stencil3d/Makefile

@@ -0,0 +1,14 @@
+ROOT_DIR := $(realpath ../../..)
+include $(ROOT_DIR)/config.mk
+
+PROJECT := stencil3d
+
+SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT)
+
+SRCS := $(SRC_DIR)/main.cpp
+
+VX_SRCS := $(SRC_DIR)/kernel.cpp
+
+OPTS ?= -n32-b2 # 32x32x32 matrix and block size of 2x2x2
+
+include ../common.mk

tests/regression/stencil3d/common.h

@@ -0,0 +1,18 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+#ifndef TYPE
+#define TYPE float
+#endif
+
+typedef struct
+{
+    uint32_t grid_dim[3];
+    uint32_t block_dim[3];
+    uint32_t size;
+    uint32_t block_size;
+    uint64_t A_addr;
+    uint64_t B_addr;
+} kernel_arg_t;
+
+#endif

tests/regression/stencil3d/kernel.cpp

@@ -0,0 +1,58 @@
+#include <vx_spawn.h>
+#include "common.h"
+
+void kernel_body(kernel_arg_t *arg)
+{
+    auto A = reinterpret_cast<TYPE *>(arg->A_addr);
+    auto B = reinterpret_cast<TYPE *>(arg->B_addr);
+    auto size = arg->size; // Assuming 'size' now represents one dimension of a cubic space.
+
+    // Calculate global column, row, and depth indices using both block and thread indices
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int dep = blockIdx.z * blockDim.z + threadIdx.z;
+
+    TYPE sum = 0;
+    int count = 0;
+
+    // Stencil kernel size is assumed to be 3x3x3
+    for (int dz = -1; dz <= 1; ++dz)
+    {
+        for (int dy = -1; dy <= 1; ++dy)
+        {
+            for (int dx = -1; dx <= 1; ++dx)
+            {
+                // Compute the neighbor's index, handling boundary conditions manually
+                int nz = dep + dz;
+                int ny = row + dy;
+                int nx = col + dx;
+
+                // Clamp the indices to be within the boundary of the array
+                if (nz < 0) {nz = 0;}
+                else if (nz >= size){
+                    nz = size - 1;}
+                if (ny < 0) {
+                    ny = 0; }
+                else if (ny >= size){
+                    ny = size - 1;}
+                if (nx < 0) {
+                    nx = 0;}
+                else if (nx >= size){
+                    nx = size - 1;}
+
+                // Add the neighbor's value to sum
+                sum += A[nz * size * size + ny * size + nx];
+                count++;
+            }
+        }
+    }
+
+    // Compute the average of the sum of neighbors and write to the output array
+    B[dep * size * size + row * size + col] = sum / count;
+}
+
+int main()
+{
+    auto arg = (kernel_arg_t *)csr_read(VX_CSR_MSCRATCH);
+    return vx_spawn_threads(3, arg->grid_dim, arg->block_dim, (vx_kernel_func_cb)kernel_body, arg);
+}