[Beignet] [PATCH 2/2] add utest for __gen_ocl_get_simd_id

[Yang, Rong R]

Pushed.

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Guo Yejun
> Sent: Tuesday, April 21, 2015 11:45
> To: beignet at lists.freedesktop.org
> Cc: Guo, Yejun
> Subject: [Beignet] [PATCH 2/2] add utest for __gen_ocl_get_simd_id
> 
> Signed-off-by: Guo Yejun <yejun.guo at intel.com>
> ---
>  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
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet