minor update

tests/opencl/common.mk

@@ -73,7 +73,7 @@ OBJS := $(addsuffix .o, $(notdir $(SRCS)))
 all: $(PROJECT) kernel.pocl
  
 kernel.pocl: kernel.cl
-	LLVM_PREFIX=$(LLVM_VORTEX) POCL_DEBUG=all LD_LIBRARY_PATH=$(LLVM_POCL)/lib:$(POCL_CC_PATH)/lib:$(LLVM_VORTEX)/lib POCL_VORTEX_CFLAGS="$(K_CFLAGS)" POCL_VORTEX_LDFLAGS="$(K_LDFLAGS)" $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
+	LD_LIBRARY_PATH=$(LLVM_POCL)/lib:$(POCL_CC_PATH)/lib:$(LLVM_VORTEX)/lib:$(LD_LIBRARY_PATH) LLVM_PREFIX=$(LLVM_VORTEX) POCL_DEBUG=all POCL_VORTEX_CFLAGS="$(K_CFLAGS)" POCL_VORTEX_LDFLAGS="$(K_LDFLAGS)" $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
  
 %.cc.o: %.cc
 	$(CXX) $(CXXFLAGS) -c $< -o $@
@@ -87,6 +87,9 @@ kernel.pocl: kernel.cl
 $(PROJECT): $(OBJS)
 	$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
 
+run-hostgpu: $(PROJECT) kernel.pocl
+	./$(PROJECT) $(OPTS)
+
 run-simx: $(PROJECT) kernel.pocl   
 	LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_RT_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
 

tests/opencl/matmul/kernel.cl

@@ -28,6 +28,9 @@ __kernel void matmul(__global float *A,
         for (int j = 0; j < localSize; j++) {
             sum += localA[localRow * localSize + j] * localB[j * localSize + localCol];
         }
+
+        // Ensure computation is done before loading next block
+        barrier(CLK_LOCAL_MEM_FENCE);
     }
 
     C[globalRow * N + globalCol] = sum;

tests/regression/tensor/main.cpp

@@ -2,6 +2,7 @@
 #include <unistd.h>
 #include <string.h>
 #include <vector>
+#include <chrono>
 #include <vortex.h>
 #include "common.h"
 
@@ -122,43 +123,6 @@ void cleanup() {
   }
 }
 
-int run_test(const kernel_arg_t& kernel_arg,
-             uint32_t buf_size, 
-             const std::vector<TYPE>& refs) {              
-  // start device
-  std::cout << "start device" << std::endl;
-  RT_CHECK(vx_start(device));
-
-  // wait for completion
-  std::cout << "wait for completion" << std::endl;
-  RT_CHECK(vx_ready_wait(device, VX_MAX_TIMEOUT));
-
-  // download destination buffer
-  std::cout << "download destination buffer" << std::endl;
-  RT_CHECK(vx_copy_from_dev(device, staging_buf.data(), kernel_arg.C_addr, buf_size));
-
-  // verify result
-  std::cout << "verify result" << std::endl;  
-  {
-    int errors = 0;
-    auto buf_ptr = (TYPE*)staging_buf.data();
-    for (uint32_t i = 0; i < refs.size(); ++i) {
-      auto ref = refs[i];
-      auto cur = buf_ptr[i];
-      if (!Comparator<TYPE>::compare(cur, ref, i, errors)) {
-        ++errors;
-      }
-    }
-    if (errors != 0) {
-      std::cout << "Found " << std::dec << errors << " errors!" << std::endl;
-      std::cout << "FAILED!" << std::endl;
-      return 1;  
-    }
-  }
-
-  return 0;
-}
-
 int main(int argc, char *argv[]) {  
   // parse command arguments
   parse_args(argc, argv);
@@ -239,10 +203,43 @@ int main(int argc, char *argv[]) {
   std::cout << "clear destination buffer" << std::endl;
   memset(staging_buf.data(), 0, num_points * sizeof(TYPE));
   RT_CHECK(vx_copy_to_dev(device, kernel_arg.C_addr, staging_buf.data(), buf_size));  
+
+  auto time_start = std::chrono::high_resolution_clock::now();
   
-  // run tests
+  // start device
-  std::cout << "run tests" << std::endl;
+  std::cout << "start device" << std::endl;
-  RT_CHECK(run_test(kernel_arg, buf_size, refs));
+  RT_CHECK(vx_start(device));
+
+  // wait for completion
+  std::cout << "wait for completion" << std::endl;
+  RT_CHECK(vx_ready_wait(device, VX_MAX_TIMEOUT));  
+
+  auto time_end = std::chrono::high_resolution_clock::now();
+  double elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(time_end - time_start).count();
+  printf("Elapsed time: %lg ms\n", elapsed);
+
+  // download destination buffer
+  std::cout << "download destination buffer" << std::endl;
+  RT_CHECK(vx_copy_from_dev(device, staging_buf.data(), kernel_arg.C_addr, buf_size));
+
+  // verify result
+  std::cout << "verify result" << std::endl;  
+  {
+    int errors = 0;
+    auto buf_ptr = (TYPE*)staging_buf.data();
+    for (uint32_t i = 0; i < refs.size(); ++i) {
+      auto ref = refs[i];
+      auto cur = buf_ptr[i];
+      if (!Comparator<TYPE>::compare(cur, ref, i, errors)) {
+        ++errors;
+      }
+    }
+    if (errors != 0) {
+      std::cout << "Found " << std::dec << errors << " errors!" << std::endl;
+      std::cout << "FAILED!" << std::endl;
+      return 1;  
+    }
+  }
 
   // cleanup
   std::cout << "cleanup" << std::endl;