This version works fine. Just pushed, thanks.
On Tue, Oct 14, 2014 at 08:08:46AM +0800, [email protected] wrote: > From: Luo Xionghu <[email protected]> > > v2: add all types to test. > v3: fix signed type count bits error. > > Signed-off-by: Luo Xionghu <[email protected]> > --- > kernels/compiler_popcount.cl | 16 +++++++++ > utests/CMakeLists.txt | 1 + > utests/compiler_popcount.cpp | 75 > ++++++++++++++++++++++++++++++++++++++++++ > 3 files changed, 92 insertions(+) > create mode 100644 kernels/compiler_popcount.cl > create mode 100644 utests/compiler_popcount.cpp > > diff --git a/kernels/compiler_popcount.cl b/kernels/compiler_popcount.cl > new file mode 100644 > index 0000000..1636118 > --- /dev/null > +++ b/kernels/compiler_popcount.cl > @@ -0,0 +1,16 @@ > +#define TEST_TYPE(TYPE) \ > +kernel void test_##TYPE(global TYPE *src, global TYPE *dst) { \ > + int i = get_global_id(0); \ > + dst[i] = popcount(src[i]); \ > +} > + > +TEST_TYPE(char) > +TEST_TYPE(uchar) > +TEST_TYPE(short) > +TEST_TYPE(ushort) > +TEST_TYPE(int) > +TEST_TYPE(uint) > +TEST_TYPE(long) > +TEST_TYPE(ulong) > + > +#undef TEST_TYPE > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index b45ecf9..1b8caca 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -41,6 +41,7 @@ set (utests_sources > compiler_ceil.cpp > compiler_clz_short.cpp > compiler_clz_int.cpp > + compiler_popcount.cpp > compiler_convert_uchar_sat.cpp > compiler_copy_buffer.cpp > compiler_copy_image.cpp > diff --git a/utests/compiler_popcount.cpp b/utests/compiler_popcount.cpp > new file mode 100644 > index 0000000..c960ae6 > --- /dev/null > +++ b/utests/compiler_popcount.cpp > @@ -0,0 +1,75 @@ > +#include "utest_helper.hpp" > + > +namespace { > + > +template<typename T> > +T get_max(); > + > +#define DEF_TEMPLATE(TYPE, NAME) \ > +template <> \ > +TYPE get_max<TYPE>() \ > +{ \ > + static TYPE max = CL_##NAME##_MAX; \ > + return max; \ > +} \ > + \ > +template <> \ > +u##TYPE get_max<u##TYPE>() \ > +{ \ > + static u##TYPE max = CL_U##NAME##_MAX; \ > + return max; \ > +} > + > +DEF_TEMPLATE(int8_t, CHAR) > +DEF_TEMPLATE(int16_t, SHRT) > +DEF_TEMPLATE(int32_t, INT) > +DEF_TEMPLATE(int64_t, LONG) > + > +template<typename T> > +void test(const char *kernel_name, int s_type) > +{ > + const int n = sizeof(T) * 8; > + > + // Setup kernel and buffers > + OCL_CREATE_KERNEL_FROM_FILE("compiler_popcount", kernel_name); > + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(T), NULL); > + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(T), NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); > + globals[0] = n; > + locals[0] = n; > + > + OCL_MAP_BUFFER(0); > + ((T*)buf_data[0])[0] = 0; > + for (int32_t i = 1; i < (int32_t) n; ++i){ > + ((T*)buf_data[0])[i] = get_max<T>() >> i; > + } > + OCL_UNMAP_BUFFER(0); > + > + OCL_NDRANGE(1); > + > + OCL_MAP_BUFFER(1); > + OCL_ASSERT(((T*)buf_data[1])[0] == 0); > + for (int i = 1; i < n; ++i){ > + OCL_ASSERT(((T*)buf_data[1])[i] == n-i-s_type); > + } > + OCL_UNMAP_BUFFER(1); > +} > + > +} > + > +#define compiler_popcount(type, kernel, s_type) \ > +static void compiler_popcount_ ##type(void)\ > +{\ > + test<type>(# kernel, s_type);\ > +}\ > +MAKE_UTEST_FROM_FUNCTION(compiler_popcount_ ## type); > + > +compiler_popcount(int8_t, test_char, 1) > +compiler_popcount(uint8_t, test_uchar, 0) > +compiler_popcount(int16_t, test_short, 1) > +compiler_popcount(uint16_t, test_ushort, 0) > +compiler_popcount(int32_t, test_int, 1) > +compiler_popcount(uint32_t, test_uint, 0) > +compiler_popcount(int64_t, test_long, 1) > +compiler_popcount(uint64_t, test_ulong, 0) > -- > 1.7.9.5 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
