diff options
| author | Luo Xionghu <xionghu.luo@intel.com> | 2014-10-27 11:14:50 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@intel.com> | 2014-10-28 10:14:46 +0800 |
| commit | 70f4cd18f52d6b7ecd1e9bae3085d1ecd57bf113 (patch) | |
| tree | 87aff4a97862b9f69bf369d58bc7101274781884 /kernels | |
| parent | 8a20f3e5e19ec13cc170861d25760e4d00a5d874 (diff) | |
add utest compiler_overflow for llvm intrinsic function.
this case only runs for uadd_with_over_flow function so far.
Signed-off-by: Luo Xionghu <xionghu.luo@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'kernels')
| -rw-r--r-- | kernels/compiler_overflow.cl | 20 |
1 files changed, 20 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) |