From: Luo Xionghu <[email protected]> only dword is supported so far.
Signed-off-by: Luo Xionghu <[email protected]> --- kernels/compiler_generic_atomic.cl | 32 +++++++++++++++++++++++++++ utests/CMakeLists.txt | 1 + utests/compiler_generic_atomic.cpp | 45 ++++++++++++++++++++++++++++++++++++++ 3 files changed, 78 insertions(+) create mode 100644 kernels/compiler_generic_atomic.cl create mode 100644 utests/compiler_generic_atomic.cpp diff --git a/kernels/compiler_generic_atomic.cl b/kernels/compiler_generic_atomic.cl new file mode 100644 index 0000000..63947ec --- /dev/null +++ b/kernels/compiler_generic_atomic.cl @@ -0,0 +1,32 @@ +#define GENERIC_KERNEL(T) \ +kernel void compiler_generic_atomic_##T(global T *src, global T *dst) \ +{ \ + size_t gid = get_global_id(0); \ + size_t lid = get_local_id(0); \ + private T pdata[16]; \ + local T ldata[16]; \ + generic T * p1 = &pdata[lid]; \ + generic T * p2 = &ldata[lid]; \ + generic T *p = (gid & 1) ? p1 : p2; \ + /* below expression is not supported by clang now */ \ + /* generic T *p = (gid & 1) ? p1 : (T *)&ldata[lid]; */ \ + *p = src[gid]; \ + /* fill other data */ \ + if(gid&1) { \ + ldata[lid] = 20; \ + } else { \ + for (int i = 0; i < 16; i++) { \ + pdata[i] = src[lid];; \ + } \ + } \ + barrier(CLK_LOCAL_MEM_FENCE); \ + \ + generic T * q1 = &pdata[lid]; \ + generic T * q2 = &ldata[lid]; \ + generic T *q = (gid & 1) ? q1 : q2; \ + dst[gid] = atomic_add(q , pdata[lid]); \ +} + +GENERIC_KERNEL(int) +//GENERIC_KERNEL(long) + diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index d1443aa..84f8fff 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -79,6 +79,7 @@ set (utests_sources compiler_global_constant_2.cpp compiler_group_size.cpp compiler_generic_pointer.cpp + compiler_generic_atomic.cpp compiler_hadd.cpp compiler_if_else.cpp compiler_integer_division.cpp diff --git a/utests/compiler_generic_atomic.cpp b/utests/compiler_generic_atomic.cpp new file mode 100644 index 0000000..9ed5f53 --- /dev/null +++ b/utests/compiler_generic_atomic.cpp @@ -0,0 +1,45 @@ +#include "utest_helper.hpp" + +template<typename T> +void test_atomic(const char* kernelName) +{ + const int n = 16; + T cpu_src[16]; + + // Setup kernel and buffers + OCL_CALL(cl_kernel_init, "compiler_generic_atomic.cl", kernelName, SOURCE, "-cl-std=CL2.0"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(T), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(T), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + + OCL_MAP_BUFFER(0); + for (int i = 0; i < n; ++i) + cpu_src[i] = ((T*)buf_data[0])[i] = (T)i; + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < n; ++i) { +// printf("i=%d dst=%d\n", i, ((T*)buf_data[1])[i]); + OCL_ASSERT(((T*)buf_data[1])[i] == 2 * cpu_src[i]); + } + OCL_UNMAP_BUFFER(1); +} + +#define GENERIC_ATOMIC_TEST(T) \ +void compiler_generic_atomic_##T() { \ + test_atomic<T>("compiler_generic_atomic_"#T); \ +} \ +MAKE_UTEST_FROM_FUNCTION(compiler_generic_atomic_##T); + +GENERIC_ATOMIC_TEST(int) +//GENERIC_TEST(long) + + + -- 2.1.4 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
