Ping for review.
-----Original Message-----
From: Beignet [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of Pan
Xiuli
Sent: Friday, December 4, 2015 1:51 PM
To: beignet@lists.freedesktop.org
Cc: Pan, Xiuli <xiuli....@intel.com>
Subject: [Beignet] [PATCH OCL2.0 V2 2/2] Utest: add a test case for built-in
ctz function
Check all type of ctz function and 0 num bound case.
Signed-off-by: Pan Xiuli <xiuli....@intel.com>
---
kernels/compiler_ctz.cl | 16 +++++++++++++
utests/CMakeLists.txt | 1 +
utests/compiler_ctz.cpp | 62 +++++++++++++++++++++++++++++++++++++++++++++++++
3 files changed, 79 insertions(+)
create mode 100644 kernels/compiler_ctz.cl create mode 100644
utests/compiler_ctz.cpp
diff --git a/kernels/compiler_ctz.cl b/kernels/compiler_ctz.cl new file mode
100644 index 0000000..8acdfb9
--- /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
6c945d2..b864465 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -116,6 +116,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 0000000..e294a03
--- /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] == i );
+ else if(sizeof(U) == 2 && i <= 16 )
+ OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ else if(sizeof(U) == 4 && i <= 32 )
+ OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ else if(sizeof(U) == 8 && i <= 64 )
+ OCL_ASSERT(((U*)buf_data[1])[i] == 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)
--
2.1.4
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
