From: Pan Xiuli <[email protected]> Check all type of ctz function and 0 num bound case.
V2: Fix type warning Signed-off-by: Pan Xiuli <[email protected]> --- 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 d1443aa..1e03697 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -117,6 +117,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..88f2109 --- /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) -- 2.5.0 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
