diff options
| -rw-r--r-- | kernels/builtin_shuffle2.cl | 13 | ||||
| -rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
| -rw-r--r-- | utests/builtin_shuffle2.cpp | 45 |
3 files changed, 59 insertions, 0 deletions
diff --git a/kernels/builtin_shuffle2.cl b/kernels/builtin_shuffle2.cl new file mode 100644 index 00000000..1a122d4a --- /dev/null +++ b/kernels/builtin_shuffle2.cl @@ -0,0 +1,13 @@ +kernel void builtin_shuffle2(global float *src1, global float *src2, global float *dst1, global float *dst2) { + int i = get_global_id(0); + float2 x = (float2)(src1[i], src2[i]); + float2 y = (float2)(1234, 5678); + uint4 mask = (uint4)(1, 0, 0, 0); + float4 v1 = shuffle2(x, y, mask); + float16 x2 = 0; + float16 y2 = (float16)(src1[i], src2[i], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0); + uint16 mask2 = (uint16)(17, 16, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0); + float16 v2 = shuffle2(x2, y2, mask2); + dst1[i] = v1.s0 + v2.s0; + dst2[i] = v1.s1 + v2.s1; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 5cd20c3f..4cdb449a 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -107,6 +107,7 @@ set (utests_sources builtin_nextafter.cpp builtin_remquo.cpp builtin_shuffle.cpp + builtin_shuffle2.cpp builtin_sign.cpp buildin_work_dim.cpp builtin_global_size.cpp diff --git a/utests/builtin_shuffle2.cpp b/utests/builtin_shuffle2.cpp new file mode 100644 index 00000000..7a9ebd15 --- /dev/null +++ b/utests/builtin_shuffle2.cpp @@ -0,0 +1,45 @@ +#include "utest_helper.hpp" + +void builtin_shuffle2(void) +{ + const int n = 32; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("builtin_shuffle2"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[3], 0, n * sizeof(float), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]); + OCL_SET_ARG(3, sizeof(cl_mem), &buf[3]); + globals[0] = n; + locals[0] = 16; + + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (int i = 0; i < n; i ++) { + ((float *)(buf_data[0]))[i] = (rand() & 15) * 0.1f; + ((float *)(buf_data[1]))[i] = (rand() & 15) * 0.1f; + } + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + + OCL_NDRANGE(1); + + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + OCL_MAP_BUFFER(2); + OCL_MAP_BUFFER(3); + for (int i = 0; i < n; i ++) { + OCL_ASSERT(2 * ((float *)(buf_data[0]))[i] == ((float *)(buf_data[3]))[i]); + OCL_ASSERT(2 * ((float *)(buf_data[1]))[i] == ((float *)(buf_data[2]))[i]); + } + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + OCL_UNMAP_BUFFER(2); + OCL_UNMAP_BUFFER(3); +} + +MAKE_UTEST_FROM_FUNCTION(builtin_shuffle2); |