[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