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

Yang Rong rong.r.yang at intel.com
Fri Jun 7 21:33:37 PDT 2013


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



More information about the Beignet mailing list