Utest: add a test case for built-in ctz function

Check all type of ctz function and 0 num bound case.

V2: Fix type warning
Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

3 files changed, 79 insertions, 0 deletions

diff --git a/kernels/compiler_ctz.cl b/kernels/compiler_ctz.cl
new file mode 100644
index 00000000..8acdfb98
--- /dev/null
+++ b/kernels/compiler_ctz.cl
@@ -0,0 +1,16 @@
+#define COMPILER_CTZ(TYPE) \
+    kernel void compiler_ctz_##TYPE(global TYPE* src, global TYPE* dst)   \
+{                                                \
+  __global TYPE* A = &src[get_global_id(0)];    \
+  __global TYPE* B = &dst[get_global_id(0)];    \
+  *B =  ctz(*A);   \
+}
+
+COMPILER_CTZ(ulong)
+COMPILER_CTZ(uint)
+COMPILER_CTZ(ushort)
+COMPILER_CTZ(uchar)
+COMPILER_CTZ(long)
+COMPILER_CTZ(int)
+COMPILER_CTZ(short)
+COMPILER_CTZ(char)
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 56add847..27451487 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -162,6 +162,7 @@ set (utests_sources
   compiler_switch.cpp
   compiler_bswap.cpp
   compiler_clz.cpp
+  compiler_ctz.cpp
   compiler_math.cpp
   compiler_atomic_functions.cpp
   compiler_async_copy.cpp
diff --git a/utests/compiler_ctz.cpp b/utests/compiler_ctz.cpp
new file mode 100644
index 00000000..88f21093
--- /dev/null
+++ b/utests/compiler_ctz.cpp
@@ -0,0 +1,62 @@
+#include "utest_helper.hpp"
+
+namespace {
+
+template<typename T>
+T get_max();
+
+template<typename U>
+void test(const char *kernel_name)
+{
+  const size_t n = 65;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_ctz", kernel_name);
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(U), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(U), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+  OCL_MAP_BUFFER(0);
+  for (uint32_t i = 0; i < n; ++i) {
+      ((U*)buf_data[0])[i] = 1l << i;
+      if(i == sizeof(U)*8)
+        ((U*)buf_data[0])[i] = 0;
+  }
+
+  OCL_UNMAP_BUFFER(0);
+
+  globals[0] = n;
+  locals[0] = 1;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i) {
+      if(sizeof(U) == 1 && i <= 8 )
+        OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
+      else if(sizeof(U) == 2 && i <= 16 )
+        OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
+      else if(sizeof(U) == 4 && i <= 32 )
+        OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
+      else if(sizeof(U) == 8 && i <= 64 )
+        OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
+    }
+  OCL_UNMAP_BUFFER(1);
+
+}
+}
+
+#define compiler_ctz(type, kernel)\
+static void compiler_ctz_ ##type(void)\
+{\
+  test<type>(# kernel);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_ctz_ ## type);
+
+compiler_ctz(uint64_t, compiler_ctz_ulong)
+compiler_ctz(uint32_t, compiler_ctz_uint)
+compiler_ctz(uint16_t, compiler_ctz_ushort)
+compiler_ctz(uint8_t, compiler_ctz_uchar)
+compiler_ctz(int64_t, compiler_ctz_long)
+compiler_ctz(int32_t, compiler_ctz_int)
+compiler_ctz(int16_t, compiler_ctz_short)
+compiler_ctz(int8_t, compiler_ctz_char)