Add tests from sgemm tutor

Author: vadyushkins

blackbox.simx.cache

@@ -0,0 +1 @@
+-DNUM_CLUSTERS=1 -DNUM_CORES=16 -DNUM_WARPS=16 -DNUM_THREADS=16    +0+0

tests/opencl/Makefile

@@ -22,6 +22,9 @@ all:
 	$(MAKE) -C sgemm2
 	$(MAKE) -C sgemm3
 	$(MAKE) -C psum
+	$(MAKE) -C kernel1
+	$(MAKE) -C kernel2
+	$(MAKE) -C kernel3
 
 run-simx:
 	$(MAKE) -C vecadd run-simx
@@ -43,6 +46,9 @@ run-simx:
 	$(MAKE) -C sgemm2 run-simx
 	$(MAKE) -C sgemm3 run-simx
 	$(MAKE) -C psum run-simx
+	$(MAKE) -C kernel1 run-simx
+	$(MAKE) -C kernel2 run-simx
+	$(MAKE) -C kernel3 run-simx
 
 run-rtlsim:
 	$(MAKE) -C vecadd run-rtlsim
@@ -64,6 +70,9 @@ run-rtlsim:
 	$(MAKE) -C sgemm2 run-rtlsim
 	$(MAKE) -C sgemm3 run-rtlsim
 	$(MAKE) -C psum run-rtlsim
+	$(MAKE) -C kernel1 run-rtlsim
+	$(MAKE) -C kernel2 run-rtlsim
+	$(MAKE) -C kernel3 run-rtlsim
 
 run-opae:
 	$(MAKE) -C vecadd run-opae
@@ -85,6 +94,9 @@ run-opae:
 	$(MAKE) -C sgemm2 run-opae
 	$(MAKE) -C sgemm3 run-opae
 	$(MAKE) -C psum run-opae
+	$(MAKE) -C kernel1 run-opae
+	$(MAKE) -C kernel2 run-opae
+	$(MAKE) -C kernel3 run-opae
 
 clean:
 	$(MAKE) -C vecadd clean
@@ -106,4 +118,7 @@ clean:
 	$(MAKE) -C bfs clean
 	$(MAKE) -C sgemm2 clean
 	$(MAKE) -C sgemm3 clean
-	$(MAKE) -C psum clean
+	$(MAKE) -C psum clean
+	$(MAKE) -C kernel1 clean
+	$(MAKE) -C kernel2 clean
+	$(MAKE) -C kernel3 clean

tests/opencl/kernel1/Makefile

@@ -0,0 +1,17 @@
+ROOT_DIR := $(realpath ../../..)
+include $(ROOT_DIR)/config.mk
+
+PROJECT := kernel1
+
+SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
+
+SRCS := $(SRC_DIR)/main.cc
+
+kernel.cl: $(SRC_DIR)/kernel.cl
+	cp $< $@
+
+KERNEL_SRCS := kernel.cl
+
+OPTS ?=
+
+include ../common.mk

tests/opencl/kernel1/kernel.cl

@@ -0,0 +1,18 @@
+__kernel void myGEMM1(const int M, const int N, const int K,
+                      const __global float* A,
+                      const __global float* B,
+                      __global float* C) {
+    // Thread identifiers
+    const int globalRow = get_global_id(0); // Row ID of C (0..M)
+    const int globalCol = get_global_id(1); // Col ID of C (0..N)
+
+    // Compute a single element (loop over K)
+    float acc = 0.0f;
+    for (int k=0; k<K; k++) {
+        acc += A[k*M + globalRow] * B[globalCol*K + k];
+    }
+
+    // Store the result
+    C[globalCol*M + globalRow] = acc;
+}
+