diff options
author | Homer Hsing <homer.xing@intel.com> | 2013-07-16 14:34:20 +0800 |
---|---|---|
committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-07-16 14:52:34 +0800 |
commit | 9373a6c91b19f3252cd6843f07d119b674262190 (patch) | |
tree | 0a5bffde7089333275c25952c6c48c09c01edab2 | |
parent | 3ae63be802a7dac3e40812dd3526f47de33fd6ce (diff) | |
download | beignet-9373a6c91b19f3252cd6843f07d119b674262190.tar.gz |
test built-in function "frexp"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
-rw-r--r-- | kernels/builtin_frexp.cl | 6 | ||||
-rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
-rw-r--r-- | utests/builtin_frexp.cpp | 50 |
3 files changed, 57 insertions, 0 deletions
diff --git a/kernels/builtin_frexp.cl b/kernels/builtin_frexp.cl new file mode 100644 index 00000000..3d33a141 --- /dev/null +++ b/kernels/builtin_frexp.cl @@ -0,0 +1,6 @@ +kernel void builtin_frexp(global float *src, global float *dst, global int *e) { + int i = get_global_id(0); + int v; + dst[i] = frexp(src[i], &v); + e[i] = v; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 14ff6b19..6f6fb863 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -101,6 +101,7 @@ set (utests_sources compiler_cl_finish.cpp get_cl_info.cpp builtin_bitselect.cpp + builtin_frexp.cpp builtin_mad_sat.cpp builtin_sign.cpp buildin_work_dim.cpp diff --git a/utests/builtin_frexp.cpp b/utests/builtin_frexp.cpp new file mode 100644 index 00000000..75dac3bc --- /dev/null +++ b/utests/builtin_frexp.cpp @@ -0,0 +1,50 @@ +#include <cmath> +#include "utest_helper.hpp" + +void builtin_frexp(void) +{ + const int n = 32; + float src[n]; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("builtin_frexp"); + 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(int), 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]); + globals[0] = n; + locals[0] = 16; + + OCL_MAP_BUFFER(0); + src[0] = ((float*)buf_data[0])[0] = 0.f; + src[1] = ((float*)buf_data[0])[1] = -0.f; + src[2] = ((float*)buf_data[0])[2] = nanf(""); + src[3] = ((float*)buf_data[0])[3] = INFINITY; + src[4] = ((float*)buf_data[0])[4] = -INFINITY; + for (int i = 5; i < n; ++i) + src[i] = ((float*)buf_data[0])[i] = (rand() & 255) * 0.1f - 12.8f; + OCL_UNMAP_BUFFER(0); + + OCL_NDRANGE(1); + + OCL_MAP_BUFFER(1); + OCL_MAP_BUFFER(2); + float *dst = (float*)buf_data[1]; + int *exp = (int*)buf_data[2]; + int w; + OCL_ASSERT(dst[0] == 0.f && exp[0] == 0); + OCL_ASSERT(dst[1] == -0.f && exp[1] == 0); + OCL_ASSERT(isnanf(dst[2])); + OCL_ASSERT(dst[3] == INFINITY); + OCL_ASSERT(dst[4] == -INFINITY); + for (int i = 5; i < n; ++i) { + OCL_ASSERT(fabsf(dst[i] - frexpf(src[i], &w)) < 1e-5); + OCL_ASSERT(exp[i] == w); + } + OCL_UNMAP_BUFFER(1); + OCL_UNMAP_BUFFER(2); +} + +MAKE_UTEST_FROM_FUNCTION(builtin_frexp); |