[Beignet] [PATCH v9 2/2] test builtin function "modf"

[Homer Hsing]

Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
 kernels/builtin_modf.cl |  6 ++++++
 utests/CMakeLists.txt   |  1 +
 utests/builtin_modf.cpp | 56 +++++++++++++++++++++++++++++++++++++++++++++++++
 3 files changed, 63 insertions(+)
 create mode 100644 kernels/builtin_modf.cl
 create mode 100644 utests/builtin_modf.cpp

diff --git a/kernels/builtin_modf.cl b/kernels/builtin_modf.cl
new file mode 100644
index 0000000..43630ed
--- /dev/null
+++ b/kernels/builtin_modf.cl
@@ -0,0 +1,6 @@
+kernel void builtin_modf(global float *src, global float *dst, global float *it) {
+  int i = get_global_id(0);
+  float x;
+  dst[i] = modf(src[i], &x);
+  it[i] = x;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index c4f4ab0..526b4a7 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -103,6 +103,7 @@ set (utests_sources
   builtin_bitselect.cpp
   builtin_frexp.cpp
   builtin_mad_sat.cpp
+  builtin_modf.cpp
   builtin_nextafter.cpp
   builtin_sign.cpp
   buildin_work_dim.cpp
diff --git a/utests/builtin_modf.cpp b/utests/builtin_modf.cpp
new file mode 100644
index 0000000..057e95e
--- /dev/null
+++ b/utests/builtin_modf.cpp
@@ -0,0 +1,56 @@
+#include <cmath>
+#include <cstring>
+#include "utest_helper.hpp"
+
+void builtin_modf(void)
+{
+  const int n = 32;
+  float src[n];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("builtin_modf");
+  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_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;
+
+  src[0] = INFINITY;
+  src[1] = -INFINITY;
+  src[2] = nanf("");
+  src[3] = 0;
+  src[4] = 1.5f;
+  src[5] = 2.5f;
+  src[6] = -2.5f;
+  src[7] = 20;
+  src[8] = 21;
+  src[9] = 89.5f;
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, n * sizeof(float));
+  OCL_UNMAP_BUFFER(0);
+
+  OCL_NDRANGE(1);
+
+  OCL_MAP_BUFFER(1);
+  OCL_MAP_BUFFER(2);
+  float *dst = (float *)buf_data[1];
+  float *it = (float *)buf_data[2];
+  OCL_ASSERT(dst[0] == 0 && it[0] == INFINITY);
+  OCL_ASSERT(dst[1] == -0.f && it[1] == -INFINITY);
+  OCL_ASSERT(isnanf(dst[2]) && isnanf(it[2]));
+  OCL_ASSERT(dst[3] == 0 && it[3] == 0);
+  OCL_ASSERT(dst[4] == 0.5f && it[4] == 1);
+  OCL_ASSERT(dst[5] == 0.5f && it[5] == 2);
+  OCL_ASSERT(dst[6] == -0.5f && it[6] == -2);
+  OCL_ASSERT(dst[7] == 0 && it[7] == 20);
+  OCL_ASSERT(dst[8] == 0 && it[8] == 21);
+  OCL_ASSERT(dst[9] == 0.5f && it[9] == 89);
+  OCL_UNMAP_BUFFER(1);
+  OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_modf);
-- 
1.8.1.2