Please ignore this patch set, I'll send a new patch set, thanks. -----Original Message----- From: Guo, Yejun Sent: Tuesday, May 12, 2015 1:34 PM To: [email protected] Cc: Guo, Yejun Subject: [PATCH 2/2] add utest for intel_sub_group_shuffle
Signed-off-by: Guo Yejun <[email protected]> --- kernels/compiler_sub_group_shuffle.cl | 15 ++++++++++++ utests/CMakeLists.txt | 3 ++- utests/compiler_sub_group_shuffle.cpp | 44 +++++++++++++++++++++++++++++++++++ 3 files changed, 61 insertions(+), 1 deletion(-) create mode 100644 kernels/compiler_sub_group_shuffle.cl create mode 100644 utests/compiler_sub_group_shuffle.cpp diff --git a/kernels/compiler_sub_group_shuffle.cl b/kernels/compiler_sub_group_shuffle.cl new file mode 100644 index 0000000..f139920 --- /dev/null +++ b/kernels/compiler_sub_group_shuffle.cl @@ -0,0 +1,15 @@ +__kernel void compiler_sub_group_shuffle(global int *dst, int c) { + int i = get_global_id(0); + if (i == 0) + dst[0] = __gen_ocl_get_simd_size(); + dst++; + + int from = i; + int o0 = __gen_ocl_get_simd_id(); + int o1 = intel_sub_group_shuffle(from, c); + int o2 = intel_sub_group_shuffle(from, 5); + dst[i*3] = o0; + dst[i*3+1] = o1; + dst[i*3+2] = o2; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index dcb3385..b68eb7b 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -210,7 +210,8 @@ set (utests_sources runtime_use_host_ptr_buffer.cpp runtime_alloc_host_ptr_buffer.cpp compiler_get_simd_size.cpp - compiler_get_simd_id.cpp) + compiler_get_simd_id.cpp + compiler_sub_group_shuffle.cpp) if (LLVM_VERSION_NODOT VERSION_GREATER 34) SET(utests_sources diff --git a/utests/compiler_sub_group_shuffle.cpp b/utests/compiler_sub_group_shuffle.cpp new file mode 100644 index 0000000..a746d75 --- /dev/null +++ b/utests/compiler_sub_group_shuffle.cpp @@ -0,0 +1,44 @@ +#include "utest_helper.hpp" + +void compiler_sub_group_shuffle(void) +{ + const size_t n = 32; + const int32_t buf_size = 3 * n + 1; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_sub_group_shuffle"); + OCL_CREATE_BUFFER(buf[0], 0, buf_size * sizeof(int), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + + int c = 3; + OCL_SET_ARG(1, sizeof(int), &c); + + globals[0] = n; + locals[0] = 16; + + OCL_MAP_BUFFER(0); + for (int32_t i = 0; i < buf_size; ++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]; + int simdsize = dst[0]; + OCL_ASSERT(simdsize == 8 || simdsize == 16); + + dst++; + for (int32_t i = 0; i < (int32_t) n; ++i){ + int round = i / simdsize; + int index = i % simdsize; + OCL_ASSERT(index == dst[3*i]); + OCL_ASSERT((round * simdsize + c) == dst[3*i+1]); + OCL_ASSERT((round * simdsize + 5) == dst[3*i+2]); + } + OCL_UNMAP_BUFFER(0); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_sub_group_shuffle); -- 1.9.1 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
