[Beignet] [PATCH] cl 2.0: add 3 function tests:to_global, to_local, to_private

[Meng Mengmeng]

cl 2.0 spec adds address space qualifier function: to_global, to_local,to_private, get_fence.
At this time, add first three function tests: to_global, to_local,to_private.
---
 kernels/builtin_to_global.cl | 48 ++++++++++++++++++++++
 utests/CMakeLists.txt        |  1 +
 utests/builtin_to_global.cpp | 94 ++++++++++++++++++++++++++++++++++++++++++++
 3 files changed, 143 insertions(+)
 create mode 100644 kernels/builtin_to_global.cl
 create mode 100644 utests/builtin_to_global.cpp

diff --git a/kernels/builtin_to_global.cl b/kernels/builtin_to_global.cl
new file mode 100644
index 0000000..8e597cb
--- /dev/null
+++ b/kernels/builtin_to_global.cl
@@ -0,0 +1,48 @@
+kernel void builtin_to_global(global int *dst, global int *src) {
+  local float lfloat;
+  lfloat = 0.0;
+  char pchar = '0';
+
+  int failures = 0;
+  if (to_global(src[get_global_id(0)]) == NULL)
+      failures++;
+  if (to_global(&lfloat) != NULL)
+     failures++;
+  if (to_global(&pchar) != NULL)
+     failures++;
+
+  dst[get_global_id(0)] = failures;
+}
+
+kernel void builtin_to_local(global int *dst, global int *src) {
+  local float lfloat;
+  lfloat = 0.0;
+  char pchar = '0';
+
+  int failures = 0;
+  if (to_local[src(get_global_id(0)]) != NULL)
+      failures++;
+  if (to_local(&lfloat) == NULL)
+     failures++;
+  if (to_local(&pchar) != NULL)
+     failures++;
+
+  dst[get_global_id(0)] = failures;
+}
+
+kernel void builtin_to_private(global int *dst, global int *src) {
+  local float lfloat;
+  lfloat = 0.0;
+  char pchar = '0';
+
+  int failures = 0;
+  if (to_private(src[get_global_id(0)]) != NULL)
+      failures++;
+  if (to_private(&lfloat) != NULL)
+     failures++;
+  if (to_private(&pchar) == NULL)
+     failures++;
+
+    dst[get_global_id(0)] = failures;
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index e7a9e26..e5c514a 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -163,6 +163,7 @@ set (utests_sources
   builtin_pow.cpp
   builtin_exp.cpp
   builtin_convert_sat.cpp
+  builtin_to_global.cpp
   sub_buffer.cpp
   runtime_createcontext.cpp
   runtime_set_kernel_arg.cpp
diff --git a/utests/builtin_to_global.cpp b/utests/builtin_to_global.cpp
new file mode 100644
index 0000000..271e48f
--- /dev/null
+++ b/utests/builtin_to_global.cpp
@@ -0,0 +1,94 @@
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+#include <string.h>
+
+#define udebug 0
+#define printf_c(...) \
+{\
+  printf("\033[1m\033[40;31m");\
+  printf( __VA_ARGS__ );\
+  printf("\033[0m");\
+}
+const int ori_data[] = {-100, -10, -1, 0, 1, 10, 100};
+const int count_input = sizeof(ori_data) / sizeof(ori_data[0]);
+int output[count_input];
+
+static void builtin_to_global(void)
+{
+  // Setup kernel and buffers
+  int i;
+
+  //OCL_CREATE_KERNEL("builtin_to_global");
+  OCL_CREATE_KERNEL("builtin_to_global");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+  globals[0] = count_input;
+  locals[0] = 1;
+  clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, sizeof(int) * count_input, ori_data, 0, NULL, NULL);
+
+  // Run the kernel
+  OCL_NDRANGE( 1 );
+  clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * count_input, output, 0, NULL, NULL);
+
+  for (i = 0; (uint)i < count_input; i++)
+      OCL_ASSERT(output[i] == 0);
+}
+
+static void builtin_to_local(void)
+{
+  // Setup kernel and buffers
+  int i;
+
+  //OCL_CREATE_KERNEL("builtin_to_global");
+  OCL_CREATE_KERNEL_FROM_FILE("builtin_to_global","builtin_to_local");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+  globals[0] = count_input;
+  locals[0] = 1;
+  clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, sizeof(int) * count_input, ori_data, 0, NULL, NULL);
+
+  // Run the kernel
+  OCL_NDRANGE( 1 );
+  clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * count_input, output, 0, NULL, NULL);
+
+  for (i = 0; (uint)i < count_input; i++)
+      OCL_ASSERT(output[i] == 0);
+}
+
+static void builtin_to_private(void)
+{
+  // Setup kernel and buffers
+  int i;
+
+  //OCL_CREATE_KERNEL("builtin_to_global");
+  OCL_CREATE_KERNEL_FROM_FILE("builtin_to_global","builtin_to_private");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, sizeof(int) * count_input, NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+  globals[0] = count_input;
+  locals[0] = 1;
+  clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, sizeof(int) * count_input, ori_data, 0, NULL, NULL);
+
+  // Run the kernel
+  OCL_NDRANGE( 1 );
+  clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * count_input, output, 0, NULL, NULL);
+
+  for (i = 0; (uint)i < count_input; i++)
+      OCL_ASSERT(output[i] == 0);
+}
+
+MAKE_UTEST_FROM_FUNCTION_KEEP_PROGRAM(builtin_to_global,true);
+MAKE_UTEST_FROM_FUNCTION_KEEP_PROGRAM(builtin_to_local, true);
+MAKE_UTEST_FROM_FUNCTION_KEEP_PROGRAM(builtin_to_private, true);
-- 
1.9.1