From: Luo Xionghu <[email protected]> this case only runs for uadd_with_over_flow function so far.
Signed-off-by: Luo Xionghu <[email protected]> --- kernels/compiler_overflow.cl | 21 ++++++++++++++ utests/CMakeLists.txt | 1 + utests/compiler_overflow.cpp | 63 ++++++++++++++++++++++++++++++++++++++++++ 3 files changed, 85 insertions(+) create mode 100644 kernels/compiler_overflow.cl create mode 100644 utests/compiler_overflow.cpp diff --git a/kernels/compiler_overflow.cl b/kernels/compiler_overflow.cl new file mode 100644 index 0000000..5874789 --- /dev/null +++ b/kernels/compiler_overflow.cl @@ -0,0 +1,21 @@ +uint add128(__global uint4 *A, uint4 B) +{ + *A += B; + uint4 carry = -convert_uint4((*A) < B); + + (*A).y += carry.x; + carry.y += ((*A).y < carry.x); + (*A).z += carry.y; + + carry.z += ((*A).z < carry.y); + (*A).w += carry.z; + return carry.w + ((*A).w < carry.z); +} + +__kernel void +compiler_overflow_uint4(__global uint4 *src, __global uint4 *dst) +{ + add128(&src[get_global_id(0)], 1); + dst[get_global_id(0)] = src[get_global_id(0)]; + +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index cf3addc..bd1c65f 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -104,6 +104,7 @@ set (utests_sources compiler_write_only.cpp compiler_write_only_shorts.cpp compiler_switch.cpp + compiler_overflow.cpp compiler_math.cpp compiler_atomic_functions.cpp compiler_async_copy.cpp diff --git a/utests/compiler_overflow.cpp b/utests/compiler_overflow.cpp new file mode 100644 index 0000000..465dbc5 --- /dev/null +++ b/utests/compiler_overflow.cpp @@ -0,0 +1,63 @@ +#include "utest_helper.hpp" + +namespace { +typedef struct { + uint32_t x; + uint32_t y; + uint32_t z; + uint32_t w; +} uint4; + +template<typename T> +void test(const char *kernel_name) +{ + const size_t n = 16; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_overflow", kernel_name); + 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]); + + OCL_MAP_BUFFER(0); + for (uint32_t i = 0; i < n; ++i) { + ((T*)buf_data[0])[i].x = 0xffffffff; + ((T*)buf_data[0])[i].y = 0xffffffff; + ((T*)buf_data[0])[i].z = 0xffffffff; + ((T*)buf_data[0])[i].w = i; + } + OCL_UNMAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < n; ++i) { + ((T*)buf_data[1])[i].x = i; + ((T*)buf_data[1])[i].y = i; + ((T*)buf_data[1])[i].z = i; + ((T*)buf_data[1])[i].w = i; + } + OCL_UNMAP_BUFFER(1); + + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < 16; ++i) { + OCL_ASSERT(((T*)buf_data[1])[i].x == 0); + OCL_ASSERT(((T*)buf_data[1])[i].y == 1); + OCL_ASSERT(((T*)buf_data[1])[i].z == 1); + OCL_ASSERT(((T*)buf_data[1])[i].w == i+2); + } + OCL_UNMAP_BUFFER(1); +} + +} + +#define compiler_overflow(type, kernel) \ +static void compiler_overflow_ ##type(void)\ +{\ + test<type>(# kernel);\ +}\ +MAKE_UTEST_FROM_FUNCTION(compiler_overflow_ ## type); + +compiler_overflow(uint4, compiler_overflow_uint4) -- 1.7.9.5 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
