[Beignet] [PATCH] New test case of ceil()

[Homer Hsing]

Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
 kernels/compiler_ceil.cl |  5 +++++
 utests/CMakeLists.txt    |  1 +
 utests/compiler_ceil.cpp | 43 +++++++++++++++++++++++++++++++++++++++++++
 3 files changed, 49 insertions(+)
 create mode 100644 kernels/compiler_ceil.cl
 create mode 100644 utests/compiler_ceil.cpp

diff --git a/kernels/compiler_ceil.cl b/kernels/compiler_ceil.cl
new file mode 100644
index 0000000..55ff425
--- /dev/null
+++ b/kernels/compiler_ceil.cl
@@ -0,0 +1,5 @@
+kernel void compiler_ceil(global float *src, global float *dst) {
+  int i = get_global_id(0);
+  dst[i] = ceil(src[i]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index b2e3c97..a02d91f 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -19,6 +19,7 @@ set (utests_sources
   compiler_array2.cpp
   compiler_array3.cpp
   compiler_byte_scatter.cpp
+  compiler_ceil.cpp
   compiler_convert_uchar_sat.cpp
   compiler_copy_buffer.cpp
   compiler_copy_image.cpp
diff --git a/utests/compiler_ceil.cpp b/utests/compiler_ceil.cpp
new file mode 100644
index 0000000..29c7551
--- /dev/null
+++ b/utests/compiler_ceil.cpp
@@ -0,0 +1,43 @@
+#include <cmath>
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, float *src, float *dst) {
+  dst[global_id] = ceilf(src[global_id]);
+}
+
+void compiler_ceil(void)
+{
+  const size_t n = 16;
+  float cpu_dst[16], cpu_src[16];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_ceil");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = 16;
+  locals[0] = 16;
+
+  // Run random tests
+  for (uint32_t pass = 0; pass < 8; ++pass) {
+    OCL_MAP_BUFFER(0);
+    for (int32_t i = 0; i < (int32_t) n; ++i)
+      cpu_src[i] = ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
+    OCL_UNMAP_BUFFER(0);
+
+    // Run the kernel on GPU
+    OCL_NDRANGE(1);
+
+    // Run on CPU
+    for (int32_t i = 0; i < (int32_t) n; ++i) cpu(i, cpu_src, cpu_dst);
+
+    // Compare
+    OCL_MAP_BUFFER(1);
+    for (int32_t i = 0; i < (int32_t) n; ++i)
+      OCL_ASSERT(((float *)buf_data[1])[i] == cpu_dst[i]);
+    OCL_UNMAP_BUFFER(1);
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_ceil);
-- 
1.8.1.2