[Beignet] [PATCH v8 2/2] test builtin function "shuffle"
Zhigang Gong
zhigang.gong at linux.intel.com
Tue Jul 23 23:40:08 PDT 2013
Homer,
All the patches before this email have been pushed to the master
branch right now.
As there is a big change in the ocl_stdlib.h due to python autogen
scripts and pch patches (merge all header files into one), you have
to manually merge the following v9/v10/... patches into current
master branch. Rebase your local branch will not help here.
The header file is move to ocl_stdlib.tmpl.h. Any question, please feel
free to contact me directly.
On Wed, Jul 24, 2013 at 11:05:11AM +0800, Homer Hsing wrote:
>
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
> kernels/builtin_shuffle.cl | 8 ++++++++
> utests/CMakeLists.txt | 1 +
> utests/builtin_shuffle.cpp | 45 +++++++++++++++++++++++++++++++++++++++++++++
> 3 files changed, 54 insertions(+)
> create mode 100644 kernels/builtin_shuffle.cl
> create mode 100644 utests/builtin_shuffle.cpp
>
> diff --git a/kernels/builtin_shuffle.cl b/kernels/builtin_shuffle.cl
> new file mode 100644
> index 0000000..ad988b9
> --- /dev/null
> +++ b/kernels/builtin_shuffle.cl
> @@ -0,0 +1,8 @@
> +kernel void builtin_shuffle(global float *src1, global float *src2, global float *dst1, global float *dst2) {
> + int i = get_global_id(0);
> + float2 src = (float2)(src1[i], src2[i]);
> + uint2 mask = (uint2)(1, 0);
> + float2 dst = shuffle(src, mask);
> + dst1[i] = dst.s0;
> + dst2[i] = dst.s1;
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 5dbae30..702b1f9 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -106,6 +106,7 @@ set (utests_sources
> builtin_modf.cpp
> builtin_nextafter.cpp
> builtin_remquo.cpp
> + builtin_shuffle.cpp
> builtin_sign.cpp
> buildin_work_dim.cpp
> builtin_global_size.cpp
> diff --git a/utests/builtin_shuffle.cpp b/utests/builtin_shuffle.cpp
> new file mode 100644
> index 0000000..c7fa86b
> --- /dev/null
> +++ b/utests/builtin_shuffle.cpp
> @@ -0,0 +1,45 @@
> +#include "utest_helper.hpp"
> +
> +void builtin_shuffle(void)
> +{
> + const int n = 32;
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("builtin_shuffle");
> + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
> + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
> + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(float), NULL);
> + OCL_CREATE_BUFFER(buf[3], 0, n * sizeof(float), NULL);
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
> + OCL_SET_ARG(3, sizeof(cl_mem), &buf[3]);
> + globals[0] = n;
> + locals[0] = 16;
> +
> + OCL_MAP_BUFFER(0);
> + OCL_MAP_BUFFER(1);
> + for (int i = 0; i < n; i ++) {
> + ((float *)(buf_data[0]))[i] = rand();
> + ((float *)(buf_data[1]))[i] = rand();
> + }
> + OCL_UNMAP_BUFFER(0);
> + OCL_UNMAP_BUFFER(1);
> +
> + OCL_NDRANGE(1);
> +
> + OCL_MAP_BUFFER(0);
> + OCL_MAP_BUFFER(1);
> + OCL_MAP_BUFFER(2);
> + OCL_MAP_BUFFER(3);
> + for (int i = 0; i < n; i ++) {
> + OCL_ASSERT(((float *)(buf_data[0]))[i] == ((float *)(buf_data[3]))[i]);
> + OCL_ASSERT(((float *)(buf_data[1]))[i] == ((float *)(buf_data[2]))[i]);
> + }
> + OCL_UNMAP_BUFFER(0);
> + OCL_UNMAP_BUFFER(1);
> + OCL_UNMAP_BUFFER(2);
> + OCL_UNMAP_BUFFER(3);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(builtin_shuffle);
> --
> 1.8.1.2
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list