diff options
author | Guo Yejun <yejun.guo@intel.com> | 2015-04-21 11:44:57 +0800 |
---|---|---|
committer | Yang Rong <rong.r.yang@intel.com> | 2015-04-24 10:31:38 +0800 |
commit | dc22a04dcb5142444f7816f1b6352e5dbef5f8c1 (patch) | |
tree | a7a2b8dfb55340c68a5fe603c108c4f73198404b | |
parent | 325e1e8e3aa2495081c77c4a129fc5743816e52d (diff) |
add utest for __gen_ocl_get_simd_id
Signed-off-by: Guo Yejun <yejun.guo@intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>
-rw-r--r-- | kernels/compiler_get_simd_id.cl | 8 | ||||
-rw-r--r-- | utests/CMakeLists.txt | 3 | ||||
-rw-r--r-- | utests/compiler_get_simd_id.cpp | 33 |
3 files changed, 43 insertions, 1 deletions
diff --git a/kernels/compiler_get_simd_id.cl b/kernels/compiler_get_simd_id.cl new file mode 100644 index 00000000..dfe625a8 --- /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 1bdb6180..dcb33850 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 00000000..ad10bf75 --- /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); |