Now clang has a new feature for block. http://reviews.llvm.org/D16928, in spec the block must be const. This will cause a Segmentation fault on llvm3.9 when bitcast the block pointer.
-----Original Message----- From: Beignet [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of Yang Rong Sent: Thursday, March 17, 2016 6:54 PM To: beignet@lists.freedesktop.org Cc: Yang, Rong R <rong.r.y...@intel.com> Subject: [Beignet] [PATCH 10/10] OCL20: add device enqueue test case. Signed-off-by: Yang Rong <rong.r.y...@intel.com> --- kernels/compiler_device_enqueue.cl | 19 +++++++++++++++++++ utests/CMakeLists.txt | 3 ++- utests/compiler_device_enqueue.cpp | 36 ++++++++++++++++++++++++++++++++++++ 3 files changed, 57 insertions(+), 1 deletion(-) create mode 100644 kernels/compiler_device_enqueue.cl create mode 100644 utests/compiler_device_enqueue.cpp diff --git a/kernels/compiler_device_enqueue.cl b/kernels/compiler_device_enqueue.cl new file mode 100644 index 0000000..6f1de7c --- /dev/null +++ b/kernels/compiler_device_enqueue.cl @@ -0,0 +1,19 @@ +void block_fn(__global uint* val) +{ + atomic_add(val, get_global_id(0)); +} + +kernel void compiler_device_enqueue(uint glob_size_arr, __global uint* +val) { + size_t tid = get_global_id(0); + + for(int i = 0; i < glob_size_arr; i++) + { + ndrange_t ndrange = ndrange_1D(glob_size_arr); + __global uint * v = val + tid; + void (^kernelBlock)(void) = ^{ block_fn(v); }; + queue_t q = get_default_queue(); + enqueue_kernel(q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, +kernelBlock); + } +} + diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index c75e6b7..ab7525f 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -227,7 +227,8 @@ set (utests_sources compiler_get_sub_group_id.cpp compiler_sub_group_shuffle.cpp runtime_pipe_query.cpp - compiler_pipe_builtin.cpp) + compiler_pipe_builtin.cpp + compiler_device_enqueue.cpp) if (LLVM_VERSION_NODOT VERSION_GREATER 34) SET(utests_sources diff --git a/utests/compiler_device_enqueue.cpp b/utests/compiler_device_enqueue.cpp new file mode 100644 index 0000000..a9e3e2d --- /dev/null +++ b/utests/compiler_device_enqueue.cpp @@ -0,0 +1,36 @@ +#include "utest_helper.hpp" + +void compiler_device_enqueue(void) +{ + const size_t n = 32; + const uint32_t global_sz = 3; + uint32_t result = 0; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_device_enqueue"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(uint32_t), &global_sz); OCL_SET_ARG(1, + sizeof(cl_mem), &buf[0]); + + OCL_MAP_BUFFER(0); + for(uint32_t i = 0; i < 69; ++i) + ((short *)buf_data[0])[i] = 0; + OCL_UNMAP_BUFFER(0); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + for(uint32_t i = 0; i < global_sz; ++i) { + result += i; + } + result *= global_sz; + + OCL_MAP_BUFFER(0); + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((uint32_t *)buf_data[0])[i] == result); + OCL_UNMAP_BUFFER(0); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_device_enqueue); -- 1.9.1 _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/beignet
