Pushed.
> -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Guo Yejun > Sent: Tuesday, April 21, 2015 11:45 > To: [email protected] > Cc: Guo, Yejun > Subject: [Beignet] [PATCH 2/2] add utest for __gen_ocl_get_simd_id > > Signed-off-by: Guo Yejun <[email protected]> > --- > kernels/compiler_get_simd_id.cl | 8 ++++++++ > utests/CMakeLists.txt | 3 ++- > utests/compiler_get_simd_id.cpp | 33 > +++++++++++++++++++++++++++++++++ > 3 files changed, 43 insertions(+), 1 deletion(-) create mode 100644 > kernels/compiler_get_simd_id.cl create mode 100644 > utests/compiler_get_simd_id.cpp > > diff --git a/kernels/compiler_get_simd_id.cl > b/kernels/compiler_get_simd_id.cl new file mode 100644 index > 0000000..dfe625a > --- /dev/null > +++ b/kernels/compiler_get_simd_id.cl > @@ -0,0 +1,8 @@ > +__kernel void compiler_get_simd_id(global int *dst) { > + int i = get_global_id(0); > + if (i == 0) > + dst[0] = __gen_ocl_get_simd_size(); > + > + dst[i+1] = __gen_ocl_get_simd_id(); > +} > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index > 1bdb618..dcb3385 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -209,7 +209,8 @@ set (utests_sources > vload_bench.cpp > runtime_use_host_ptr_buffer.cpp > runtime_alloc_host_ptr_buffer.cpp > - compiler_get_simd_size.cpp) > + compiler_get_simd_size.cpp > + compiler_get_simd_id.cpp) > > if (LLVM_VERSION_NODOT VERSION_GREATER 34) > SET(utests_sources > diff --git a/utests/compiler_get_simd_id.cpp > b/utests/compiler_get_simd_id.cpp new file mode 100644 index > 0000000..ad10bf7 > --- /dev/null > +++ b/utests/compiler_get_simd_id.cpp > @@ -0,0 +1,33 @@ > +#include "utest_helper.hpp" > + > +void compiler_get_simd_id(void) > +{ > + const size_t n = 256; > + > + // Setup kernel and buffers > + OCL_CREATE_KERNEL("compiler_get_simd_id"); > + OCL_CREATE_BUFFER(buf[0], 0, (n+1) * sizeof(int), NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + > + globals[0] = n; > + locals[0] = 16; > + > + OCL_MAP_BUFFER(0); > + for (int32_t i = 0; i < (int32_t) (n+1); ++i) > + ((int*)buf_data[0])[i] = -1; > + OCL_UNMAP_BUFFER(0); > + > + // Run the kernel on GPU > + OCL_NDRANGE(1); > + > + // Compare > + OCL_MAP_BUFFER(0); > + int* dst = (int *)buf_data[0]; > + OCL_ASSERT(8 == dst[0] || 16 == dst[0]); > + for (int32_t i = 1; i < (int32_t) n; ++i){ > + OCL_ASSERT((i-1) % dst[0] == dst[i]); > + } > + OCL_UNMAP_BUFFER(0); > +} > + > +MAKE_UTEST_FROM_FUNCTION(compiler_get_simd_id); > -- > 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
