[Beignet] [PATCH 2/2] Fix two tests fail when OCL_SIMD_WIDTH=8.

[Zhigang Gong]

Pushed, thanks for the patchset.

On Sat, Jun 08, 2013 at 12:33:37PM +0800, Yang Rong wrote:
> Add barrier for compiler_local_memory and compiler_local_memory_two_ptr,
> otherwise tests may fail if work group size bigger than thread's simd size.
> 
> After add barrier, the test compiler_local_memory is same as
> compiler_local_memory_barrier, so delete test compiler_local_memory.
> 
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
>  kernels/compiler_local_memory.cl         |    5 ----
>  kernels/compiler_local_memory_two_ptr.cl |    1 +
>  utests/CMakeLists.txt                    |    1 -
>  utests/compiler_local_memory.cpp         |   47 ------------------------------
>  4 files changed, 1 insertion(+), 53 deletions(-)
>  delete mode 100644 kernels/compiler_local_memory.cl
>  delete mode 100644 utests/compiler_local_memory.cpp
> 
> diff --git a/kernels/compiler_local_memory.cl b/kernels/compiler_local_memory.cl
> deleted file mode 100644
> index daadd66..0000000
> --- a/kernels/compiler_local_memory.cl
> +++ /dev/null
> @@ -1,5 +0,0 @@
> -__kernel void compiler_local_memory(__global int *dst, __local int *src) {
> -  src[get_local_id(0)] = get_local_id(0);
> -  dst[get_global_id(0)] = src[15 - get_local_id(0)];
> -}
> -
> diff --git a/kernels/compiler_local_memory_two_ptr.cl b/kernels/compiler_local_memory_two_ptr.cl
> index f410406..46589ba 100644
> --- a/kernels/compiler_local_memory_two_ptr.cl
> +++ b/kernels/compiler_local_memory_two_ptr.cl
> @@ -4,6 +4,7 @@ __kernel void compiler_local_memory_two_ptr(__global int *dst,
>  {
>    src0[get_local_id(0)] = get_local_id(0);
>    src1[get_local_id(0)] = get_global_id(0);
> +  barrier(CLK_LOCAL_MEM_FENCE);
>    dst[get_global_id(0)] = src0[15 - get_local_id(0)] + src1[15 - get_local_id(0)];
>  }
>  
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 56685b0..d883cfd 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -71,7 +71,6 @@ set (utests_sources
>    compiler_insn_selection_min.cpp
>    compiler_insn_selection_max.cpp
>    compiler_insn_selection_masked_min_max.cpp
> -  compiler_local_memory.cpp
>    compiler_local_memory_two_ptr.cpp
>    compiler_local_memory_barrier.cpp
>    compiler_local_memory_barrier_wg64.cpp
> diff --git a/utests/compiler_local_memory.cpp b/utests/compiler_local_memory.cpp
> deleted file mode 100644
> index 49fa28c..0000000
> --- a/utests/compiler_local_memory.cpp
> +++ /dev/null
> @@ -1,47 +0,0 @@
> -/* 
> - * Copyright © 2012 Intel Corporation
> - *
> - * This library is free software; you can redistribute it and/or
> - * modify it under the terms of the GNU Lesser General Public
> - * License as published by the Free Software Foundation; either
> - * version 2 of the License, or (at your option) any later version.
> - *
> - * This library is distributed in the hope that it will be useful,
> - * but WITHOUT ANY WARRANTY; without even the implied warranty of
> - * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
> - * Lesser General Public License for more details.
> - *
> - * You should have received a copy of the GNU Lesser General Public
> - * License along with this library. If not, see <http://www.gnu.org/licenses/>.
> - *
> - * Author: Benjamin Segovia <benjamin.segovia at intel.com>
> - */
> -
> -#include "utest_helper.hpp"
> -
> -static void compiler_local_memory(void)
> -{
> -  const size_t n = 1024;
> -
> -  // Setup kernel and buffers
> -  OCL_CREATE_KERNEL("compiler_local_memory");
> -  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
> -  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> -  OCL_SET_ARG(1, 64, NULL); // 16 x int
> -
> -  // Run the kernel
> -  globals[0] = n;
> -  locals[0] = 16;
> -  OCL_NDRANGE(1);
> -  OCL_MAP_BUFFER(0);
> -
> -  // Check results
> -  uint32_t *dst = (uint32_t*)buf_data[0];
> -  for (uint32_t i = 0; i < n; i+=16)
> -  for (uint32_t j = 0; j < 16; ++j)
> -    OCL_ASSERT(dst[i+j] == 15-j);
> -}
> -
> -MAKE_UTEST_FROM_FUNCTION(compiler_local_memory);
> -
> -
> -- 
> 1.7.10.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet