Zero is definitely covered as the thread number is 64, all maximum of type value will be zero after shifting the size of type bites. I've tested the signed type LZD kernel, but not easy to include it in this utest, should add it in another patch.
Luo Xionghu Best Regards -----Original Message----- From: Song, Ruiling Sent: Thursday, January 15, 2015 4:00 PM To: Luo, Xionghu; [email protected] Cc: Luo, Xionghu Subject: RE: [Beignet] [PATCH 2/3] add clz(count leading zero) utest. __builtin_clz only accept positive number? Why do you only cover positive number? I think negative number and zero need also be covered. > -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf > Of [email protected] > Sent: Thursday, January 15, 2015 1:22 PM > To: [email protected] > Cc: Luo, Xionghu > Subject: [Beignet] [PATCH 2/3] add clz(count leading zero) utest. > > From: Luo Xionghu <[email protected]> > > this kernl calls the llvm __builtin_clz to generate llvm.clz function > then call the gen instruction clz, different from the test > compiler_clz_int, which use the fbh to implement. > > Signed-off-by: Luo Xionghu <[email protected]> > --- > kernels/compiler_clz.cl | 12 +++++++++ > utests/CMakeLists.txt | 1 + > utests/compiler_clz.cpp | 67 > +++++++++++++++++++++++++++++++++++++++++++++++++ > 3 files changed, 80 insertions(+) > create mode 100644 kernels/compiler_clz.cl create mode 100644 > utests/compiler_clz.cpp > > diff --git a/kernels/compiler_clz.cl b/kernels/compiler_clz.cl new > file mode > 100644 index 0000000..7ab6261 > --- /dev/null > +++ b/kernels/compiler_clz.cl > @@ -0,0 +1,12 @@ > +#define COMPILER_CLZ(TYPE) \ > + kernel void compiler_clz_##TYPE(global TYPE* src, global TYPE* > +dst) > \ > +{ \ > + __global TYPE* A = &src[get_global_id(0)]; \ > + __global TYPE* B = &dst[get_global_id(0)]; \ > + *B = __builtin_clz(*A); \ > +} > + > +COMPILER_CLZ(uint) > +COMPILER_CLZ(ulong) > +COMPILER_CLZ(ushort) > +COMPILER_CLZ(uchar) > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index > 5b29c0b..193fef3 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -105,6 +105,7 @@ set (utests_sources > compiler_write_only_shorts.cpp > compiler_switch.cpp > compiler_bswap.cpp > + compiler_clz.cpp > compiler_math.cpp > compiler_atomic_functions.cpp > compiler_async_copy.cpp > diff --git a/utests/compiler_clz.cpp b/utests/compiler_clz.cpp new > file mode > 100644 index 0000000..901e19b > --- /dev/null > +++ b/utests/compiler_clz.cpp > @@ -0,0 +1,67 @@ > +#include "utest_helper.hpp" > + > +namespace { > + > +template <typename U> > +U get_max() > +{ > + int shift_bit = sizeof(U)*8; > + U u_max = 0; > + for (int i = 0; i < shift_bit; i++) > + u_max |= 1<<(shift_bit-i-1); > + return u_max; > +} > + > +template<typename U> > +void test(const char *kernel_name) > +{ > + const size_t n = 64; > + > + // Setup kernel and buffers > + OCL_CREATE_KERNEL_FROM_FILE("compiler_clz", 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]); > + > + U max = get_max<U>(); > + > + OCL_MAP_BUFFER(0); > + for (uint32_t i = 0; i < n; ++i) { > + ((U*)buf_data[0])[i] = max >> i; } OCL_UNMAP_BUFFER(0); > + > + globals[0] = n; > + locals[0] = 16; > + 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+24) ); > + else if(sizeof(U) == 2 && i < 16 ) > + OCL_ASSERT(((U*)buf_data[1])[i] == (i+16) ); > + else if(sizeof(U) == 4 && i < 32 ) > + OCL_ASSERT(((U*)buf_data[1])[i] == i ); > + else if(sizeof(U) == 8 && i < 32 ) > + OCL_ASSERT(((U*)buf_data[1])[i] == 0 ); > + else if(sizeof(U) == 8 && i > 31) > + OCL_ASSERT(((U*)buf_data[1])[i] == (i-32) ); } > + OCL_UNMAP_BUFFER(1); > + > +} > + > +} > + > +#define compiler_clz(type, kernel) \ > +static void compiler_clz_ ##type(void)\ {\ > + test<type>(# kernel);\ > +}\ > +MAKE_UTEST_FROM_FUNCTION(compiler_clz_ ## type); > + > +compiler_clz(uint64_t, compiler_clz_ulong) compiler_clz(uint32_t, > +compiler_clz_uint) compiler_clz(uint16_t, compiler_clz_ushort) > +compiler_clz(uint8_t, compiler_clz_uchar) > -- > 1.9.1 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
