[Beignet] [PATCH v3 4/4] Add constant ptr argument test case.

[Yang Rong]

Signed-off-by: Yang Rong <rong.r.yang at intel.com>
---
 kernels/compiler_function_constant.cl  |    6 ++++
 kernels/compiler_function_constant0.cl |    6 ++++
 utests/CMakeLists.txt                  |    3 ++
 utests/compiler_function_constant.cpp  |   34 +++++++++++++++++++++++
 utests/compiler_function_constant0.cpp |   42 ++++++++++++++++++++++++++++
 utests/compiler_function_constant1.cpp |   47 ++++++++++++++++++++++++++++++++
 6 files changed, 138 insertions(+)
 create mode 100644 kernels/compiler_function_constant.cl
 create mode 100644 kernels/compiler_function_constant0.cl
 create mode 100644 utests/compiler_function_constant.cpp
 create mode 100644 utests/compiler_function_constant0.cpp
 create mode 100644 utests/compiler_function_constant1.cpp

diff --git a/kernels/compiler_function_constant.cl b/kernels/compiler_function_constant.cl
new file mode 100644
index 0000000..ca7e874
--- /dev/null
+++ b/kernels/compiler_function_constant.cl
@@ -0,0 +1,6 @@
+__kernel void
+compiler_function_constant(__constant short *c, __global int *dst, int value)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = value + c[id%69];
+}
diff --git a/kernels/compiler_function_constant0.cl b/kernels/compiler_function_constant0.cl
new file mode 100644
index 0000000..f6efcef
--- /dev/null
+++ b/kernels/compiler_function_constant0.cl
@@ -0,0 +1,6 @@
+__kernel void
+compiler_function_constant0(__constant short *c0, __constant char *c1, __global int *dst, int value)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = value + c0[id%69] + c1[15];
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index bed0159..b2e3c97 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -29,6 +29,9 @@ set (utests_sources
   compiler_function_argument0.cpp
   compiler_function_argument1.cpp
   compiler_function_argument.cpp
+  compiler_function_constant0.cpp
+  compiler_function_constant1.cpp
+  compiler_function_constant.cpp
   compiler_if_else.cpp
   compiler_integer_division.cpp
   compiler_integer_remainder.cpp
diff --git a/utests/compiler_function_constant.cpp b/utests/compiler_function_constant.cpp
new file mode 100644
index 0000000..20f0ece
--- /dev/null
+++ b/utests/compiler_function_constant.cpp
@@ -0,0 +1,34 @@
+#include "utest_helper.hpp"
+
+void compiler_function_constant(void)
+{
+  const size_t n = 2048;
+  const uint32_t value = 34;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_constant");
+  OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(uint32_t), &value);
+
+  OCL_MAP_BUFFER(0);
+  for(uint32_t i = 0; i < 69; ++i)
+    ((short *)buf_data[0])[i] = i;
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(1);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((uint32_t *)buf_data[1])[i] == (value + i%69));
+
+  OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_constant);
diff --git a/utests/compiler_function_constant0.cpp b/utests/compiler_function_constant0.cpp
new file mode 100644
index 0000000..de564f3
--- /dev/null
+++ b/utests/compiler_function_constant0.cpp
@@ -0,0 +1,42 @@
+#include "utest_helper.hpp"
+
+void compiler_function_constant0(void)
+{
+  const size_t n = 2048;
+  const uint32_t value = 34;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_constant0");
+  OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, 256 * sizeof(char), NULL);
+  OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+  OCL_SET_ARG(3, sizeof(uint32_t), &value);
+
+  OCL_MAP_BUFFER(0);
+  for(uint32_t i = 0; i < 69; ++i)
+    ((short *)buf_data[0])[i] = i;
+  OCL_UNMAP_BUFFER(0);
+
+  OCL_MAP_BUFFER(1);
+  for(uint32_t i = 0; i < 256; ++i)
+    ((char *)buf_data[1])[i] = 10;
+  ((char *)buf_data[1])[15] = 15;
+  OCL_UNMAP_BUFFER(1);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(2);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((uint32_t *)buf_data[2])[i] == (value + 15 + i%69));
+
+  OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_constant0);
diff --git a/utests/compiler_function_constant1.cpp b/utests/compiler_function_constant1.cpp
new file mode 100644
index 0000000..b92e6ca
--- /dev/null
+++ b/utests/compiler_function_constant1.cpp
@@ -0,0 +1,47 @@
+#include "utest_helper.hpp"
+
+void compiler_function_constant1(void)
+{
+  const size_t n = 2048;
+  const uint32_t value = 34;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_function_constant");
+  OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(uint32_t), &value);
+
+  OCL_MAP_BUFFER(0);
+  for(uint32_t i = 0; i < 69; ++i)
+    ((short *)buf_data[0])[i] = i;
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  OCL_CREATE_BUFFER(buf[2], 0, 101 * sizeof(short), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[2]);
+  OCL_MAP_BUFFER(2);
+  for(uint32_t i = 0; i < 69; ++i)
+    ((short *)buf_data[2])[i] = 2*i;
+  OCL_UNMAP_BUFFER(2);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  OCL_MAP_BUFFER(1);
+
+  // Check results
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((uint32_t *)buf_data[1])[i] == (value + (i%69)*2));
+
+  OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_constant1);
-- 
1.7.9.5