utests: Added one test case for clFinish().

Reviewed-by: Zou, Nanhai <nanhai.zou@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

diff --git a/kernels/test_cl_finish.cl b/kernels/test_cl_finish.cl
new file mode 100644 (file)
index 0000000..723949c
--- /dev/null
+++ b/kernels/test_cl_finish.cl
@@ -0,0 +1,12 @@
+
+
+__kernel void
+test_cl_finish(__global int *src, __global int *dst, int n, int num_threads)
+{
+       int tid, pos;
+
+       tid = get_global_id(0);
+       for (pos=tid; pos < n; pos+=num_threads) {
+               dst[pos] = src[pos];
+       }
+}

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index d883cfd..81ff5d8 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -79,6 +79,7 @@ set (utests_sources
   compiler_copy_image1.cpp
   compiler_get_image_info.cpp
   compiler_vector_load_store.cpp
+  compiler_cl_finish.cpp
   buildin_work_dim.cpp
   runtime_createcontext.cpp
   utest_assert.cpp

diff --git a/utests/compiler_cl_finish.cpp b/utests/compiler_cl_finish.cpp
new file mode 100644 (file)
index 0000000..7c7dee3
--- /dev/null
+++ b/utests/compiler_cl_finish.cpp
@@ -0,0 +1,50 @@
+#include "utest_helper.hpp"
+#include <sys/time.h>
+
+#define T_GET(t)        gettimeofday(&t, NULL);
+#define T_LAPSE(t1, t2) \
+  ((t2.tv_sec+t2.tv_usec*0.000001) - (t1.tv_sec+t1.tv_usec*0.000001))
+
+static void compiler_cl_finish(void)
+{
+  const size_t n = 16*1024*1024;
+  struct timeval t1, t2;
+  float t_fin, t_map_w_fin,t_map_wo_fin;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("test_cl_finish");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);
+
+  // Run the kernel
+  locals[0]  = 64;
+  globals[0] = 32 * locals[0];
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(int), &n);
+  OCL_SET_ARG(3, sizeof(int), &globals[0]);
+
+  // 1st time map after clFinish
+  OCL_NDRANGE(1);
+  T_GET(t1);
+  OCL_FINISH();
+  T_GET(t2);
+  t_fin = T_LAPSE(t1, t2);
+
+  T_GET(t1);
+  OCL_MAP_BUFFER(0);
+  T_GET(t2);
+  t_map_w_fin = T_LAPSE(t1, t2);
+
+  // 2nd time map without clFinish
+  OCL_NDRANGE(1);
+  T_GET(t1);
+  OCL_MAP_BUFFER(0);
+  T_GET(t2);
+  t_map_wo_fin = T_LAPSE(t1, t2);
+
+  OCL_ASSERT(t_fin > t_map_w_fin && t_map_wo_fin > t_map_w_fin);
+  OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_cl_finish);

diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index c28615e..dd98a1f 100644 (file)
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -75,6 +75,11 @@ extern EGLSurface  eglSurface;
     OCL_CALL(clFlush, queue); \
   } while(0)
 
+#define OCL_FINISH() \
+  do { \
+    OCL_CALL(clFinish, queue); \
+  } while(0)
+
 #define OCL_CALL2(FN, RET, ...) \
   do { \
     cl_int status; \