diff options
| -rw-r--r-- | kernels/compiler_overflow.cl | 20 | ||||
| -rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
| -rw-r--r-- | utests/compiler_overflow.cpp | 72 |
3 files changed, 93 insertions, 0 deletions
diff --git a/kernels/compiler_overflow.cl b/kernels/compiler_overflow.cl new file mode 100644 index 00000000..75ed5ce2 --- /dev/null +++ b/kernels/compiler_overflow.cl @@ -0,0 +1,20 @@ +#define COMPILER_OVERFLOW(TYPE) \ + kernel void compiler_overflow_##TYPE (global TYPE* src, global TYPE* dst) \ +{ \ + __global TYPE* A = &src[get_global_id(0)]; \ + TYPE B = 1; \ + *A += B; \ + TYPE carry = -convert_##TYPE((*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; \ + dst[get_global_id(0)] = src[get_global_id(0)]; \ +} + +COMPILER_OVERFLOW(uint4) +COMPILER_OVERFLOW(ushort4) +COMPILER_OVERFLOW(uchar4) diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index cf3addcb..bd1c65fe 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 00000000..1d3f53d1 --- /dev/null +++ b/utests/compiler_overflow.cpp @@ -0,0 +1,72 @@ +#include "utest_helper.hpp" + +namespace { + +typedef struct { + uint32_t x; + uint32_t y; + uint32_t z; + uint32_t w; +} uint4; + +typedef struct { + uint16_t x; + uint16_t y; + uint16_t z; + uint16_t w; +} ushort4; + +typedef struct { + uint8_t x; + uint8_t y; + uint8_t z; + uint8_t w; +} uchar4; + +template<typename T> +void test(const char *kernel_name, int s_type) +{ + 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 = s_type?CL_INT_MAX:CL_UINT_MAX; + ((T*)buf_data[0])[i].y = s_type?CL_INT_MAX:CL_UINT_MAX; + ((T*)buf_data[0])[i].z = s_type?CL_INT_MAX:CL_UINT_MAX; + ((T*)buf_data[0])[i].w = i; + } + OCL_UNMAP_BUFFER(0); + + 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, s_type) \ +static void compiler_overflow_ ##type(void)\ +{\ + test<type>(# kernel, s_type);\ +}\ +MAKE_UTEST_FROM_FUNCTION(compiler_overflow_ ## type); + +compiler_overflow(uint4, compiler_overflow_uint4, 0) +compiler_overflow(ushort4, compiler_overflow_ushort4, 0) +compiler_overflow(uchar4, compiler_overflow_uchar4, 0) |