summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGuo Yejun <yejun.guo@intel.com>2015-04-16 08:51:06 +0800
committerYang Rong <rong.r.yang@intel.com>2015-04-24 10:31:29 +0800
commit5c6220a21dc64c2165c81ae499cf7236093da1b2 (patch)
tree2df2830827fdf18152caeb69a7b8748cc44aa6bb
parente2019e40f0fe73acaa057689592accff9553e0be (diff)
add utest for __gen_ocl_get_simd_size
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_size.cl5
-rw-r--r--utests/CMakeLists.txt3
-rw-r--r--utests/compiler_get_simd_size.cpp32
3 files changed, 39 insertions, 1 deletions
diff --git a/kernels/compiler_get_simd_size.cl b/kernels/compiler_get_simd_size.cl
new file mode 100644
index 00000000..6e303a35
--- /dev/null
+++ b/kernels/compiler_get_simd_size.cl
@@ -0,0 +1,5 @@
+__kernel void compiler_get_simd_size(global int *dst)
+{
+ int i = get_global_id(0);
+ dst[i] = __gen_ocl_get_simd_size();
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 06baa689..1bdb6180 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -208,7 +208,8 @@ set (utests_sources
compiler_assignment_operation_in_if.cpp
vload_bench.cpp
runtime_use_host_ptr_buffer.cpp
- runtime_alloc_host_ptr_buffer.cpp)
+ runtime_alloc_host_ptr_buffer.cpp
+ compiler_get_simd_size.cpp)
if (LLVM_VERSION_NODOT VERSION_GREATER 34)
SET(utests_sources
diff --git a/utests/compiler_get_simd_size.cpp b/utests/compiler_get_simd_size.cpp
new file mode 100644
index 00000000..ea70cd96
--- /dev/null
+++ b/utests/compiler_get_simd_size.cpp
@@ -0,0 +1,32 @@
+#include "utest_helper.hpp"
+
+void compiler_get_simd_size(void)
+{
+ const size_t n = 256;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_get_simd_size");
+ OCL_CREATE_BUFFER(buf[0], 0, n * 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; ++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];
+ for (int32_t i = 0; i < (int32_t) n; ++i){
+ OCL_ASSERT(8 == dst[i] || 16 == dst[i]);
+ }
+ OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_get_simd_size);