[Beignet] [PATCH v2 8/8] [OCL20] utest: add atomic opencl-2.0 case to test api.

[xionghu.luo at intel.com]

From: Luo Xionghu <xionghu.luo at intel.com>

Signed-off-by: Luo Xionghu <xionghu.luo at intel.com>
---
 kernels/compiler_atomic_functions.cl | 54 ++++++++++++++++++++++++++++++++++++
 utests/compiler_atomic_functions.cpp | 16 +++++++++--
 2 files changed, 67 insertions(+), 3 deletions(-)

diff --git a/kernels/compiler_atomic_functions.cl b/kernels/compiler_atomic_functions.cl
index fbc16fb..4af17f1 100644
--- a/kernels/compiler_atomic_functions.cl
+++ b/kernels/compiler_atomic_functions.cl
@@ -48,3 +48,57 @@ __kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __g
       atomic_xchg(&dst[i+12], tmp[i]);
   }
 }
+
+__kernel void compiler_atomic_functions_20(__global int *dst, __local int *tmp, __global int *src) {
+  int lid = get_local_id(0);
+  int i = lid % 12;
+  atomic_int* p = (atomic_int*)tmp;
+  if(lid == 0) {
+    for(int j=0; j<12; j=j+1) {
+      atomic_exchange(&p[j], 0);
+    }
+    atomic_exchange(&p[4], -1);
+  }
+  barrier(CLK_LOCAL_MEM_FENCE);
+  int compare = 0;
+
+  switch(i) {
+    case 0: atomic_inc(&tmp[i]); break;
+    case 1: atomic_dec(&tmp[i]); break;
+    case 2: atomic_fetch_add(&p[i], src[lid]); break;
+    case 3: atomic_fetch_sub(&p[i], src[lid]); break;
+    case 4: atomic_fetch_and(&p[i], ~(src[lid]<<(lid / 16))); break;
+    case 5: atomic_fetch_or (&p[i], src[lid]<<(lid / 16)); break;
+    case 6: atomic_fetch_xor(&p[i], src[lid]); break;
+    case 7: atomic_fetch_min(&p[i], -src[lid]); break;
+    case 8: atomic_fetch_max(&p[i], src[lid]); break;
+    case 9: atomic_fetch_min((atomic_uint*)&p[i], -src[lid]); break;
+    case 10: atomic_fetch_max((atomic_uint*)&p[i], src[lid]); break;
+    case 11: atomic_compare_exchange_strong(&p[i], &compare, src[10]); break;
+    default:  break;
+  }
+
+  atomic_int* d = (atomic_int*)dst;
+  switch(i) {
+    case 0: atomic_inc(&dst[i]); break;
+    case 1: atomic_dec(&dst[i]); break;
+    case 2: atomic_fetch_add(&d[i], src[lid]); break;
+    case 3: atomic_fetch_sub(&d[i], src[lid]); break;
+    case 4: atomic_fetch_and(&d[i], ~(src[lid]<<(lid / 16))); break;
+    case 5: atomic_fetch_or (&d[i], src[lid]<<(lid / 16)); break;
+    case 6: atomic_fetch_xor(&d[i], src[lid]); break;
+    case 7: atomic_fetch_min(&d[i], -src[lid]); break;
+    case 8: atomic_fetch_max(&d[i], src[lid]); break;
+    case 9: atomic_fetch_min((atomic_uint*)&d[i], -src[lid]); break;
+    case 10: atomic_fetch_max((atomic_uint*)&d[i], src[lid]); break;
+    case 11: atomic_compare_exchange_strong(&d[i], &compare, src[10]); break;
+    default:  break;
+  }
+
+  barrier(CLK_GLOBAL_MEM_FENCE);
+
+  if(get_global_id(0) == 0) {
+    for(i=0; i<12; i=i+1)
+      atomic_xchg(&dst[i+12], tmp[i]);
+  }
+}
diff --git a/utests/compiler_atomic_functions.cpp b/utests/compiler_atomic_functions.cpp
index 65f1c5a..7a298cf 100644
--- a/utests/compiler_atomic_functions.cpp
+++ b/utests/compiler_atomic_functions.cpp
@@ -56,7 +56,7 @@ static void cpu_compiler_atomic(int *dst, int *src)
     dst[i+12] = tmp[i];
 }
 
-static void compiler_atomic_functions(void)
+static void compiler_atomic_functions(const char* kernel_name)
 {
   const size_t n = GROUP_NUM * LOCAL_SIZE;
   int cpu_dst[24] = {0}, cpu_src[256];
@@ -65,7 +65,7 @@ static void compiler_atomic_functions(void)
   locals[0] = LOCAL_SIZE;
 
   // Setup kernel and buffers
-  OCL_CREATE_KERNEL("compiler_atomic_functions");
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_atomic_functions", kernel_name);
   OCL_CREATE_BUFFER(buf[0], 0, 24 * sizeof(int), NULL);
   OCL_CREATE_BUFFER(buf[1], 0, locals[0] * sizeof(int), NULL);
   OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
@@ -94,4 +94,14 @@ static void compiler_atomic_functions(void)
   OCL_UNMAP_BUFFER(0);
 }
 
-MAKE_UTEST_FROM_FUNCTION(compiler_atomic_functions)
+#define compiler_atomic(kernel, version) \
+static void compiler_atomic_functions_##version()\
+{\
+  compiler_atomic_functions(kernel); \
+} \
+MAKE_UTEST_FROM_FUNCTION(compiler_atomic_functions_##version)
+
+compiler_atomic("compiler_atomic_functions", 12)
+compiler_atomic("compiler_atomic_functions_20", 20)
+
+
-- 
2.1.4