[Beignet] [PATCH] utests: add utest to convert_float_roundmode(double x)

[rander]

Signed-off-by: rander <rander.wang at intel.com>
---
 kernels/builtin_convert_double2float.cl |  25 ++++++++
 utests/CMakeLists.txt                   |   3 +-
 utests/builtin_convert_double2float.cpp | 101 ++++++++++++++++++++++++++++++++
 3 files changed, 128 insertions(+), 1 deletion(-)
 create mode 100644 kernels/builtin_convert_double2float.cl
 create mode 100644 utests/builtin_convert_double2float.cpp

diff --git a/kernels/builtin_convert_double2float.cl b/kernels/builtin_convert_double2float.cl
new file mode 100644
index 0000000..dccec03
--- /dev/null
+++ b/kernels/builtin_convert_double2float.cl
@@ -0,0 +1,25 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_convert_double2float(__global double *X,
+												__global float *Z,
+												int max_input)
+{
+	int i = get_global_id(0);
+	int j;
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_float(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_float_rtz(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_float_rtn(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_float_rte(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_float_rtp(X[j]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 6e41eeb..d76bba7 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -305,7 +305,8 @@ set (utests_sources
   builtin_convert_double2int32.cpp
   builtin_convert_double2int64.cpp
   builtin_convert_int8toDouble.cpp
-  builtin_convert_int16toDouble.cpp)
+  builtin_convert_int16toDouble.cpp
+  builtin_convert_double2float.cpp)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/builtin_convert_double2float.cpp b/utests/builtin_convert_double2float.cpp
new file mode 100644
index 0000000..94d3b6e
--- /dev/null
+++ b/utests/builtin_convert_double2float.cpp
@@ -0,0 +1,101 @@
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+namespace{
+    unsigned long doubleHX[] = {
+        0x0,
+        0x8000000000000000,  //-0
+        0x0000000000000001,  // denorm
+        0xFFFFFFFFFFFFFFFF,  //nan
+        0xFFF0000000000000,  // -inf
+        0x7FF0000000000000, //+inf
+};
+
+double doubleX[] = {
+    0x1.0p128,
+    0x1.01p128,
+    0x1.01p127,
+    0x1.fffffffffp127,
+    -0x1.0p126,
+    0x1.000001p127,
+    0x1.7p127,
+    0x1.7ffffffffp127,
+    0x1.1fffffffffp64,
+    -0x1.1fffffffffp40,
+    -0x1.01p127,
+    0x1.0p-128,
+    -0x1.7p126,
+    -0x1.0p-64,
+};
+
+const char*  testFunc[] =
+{
+    " char convert_float(double x)",
+    " char convert_float_sat(double x)",
+    " char convert_float_rtz(double x)",
+    " char convert_float_rtn(double x)",
+    " char convert_float_rte(double x)",
+};
+
+unsigned int expectResult[] = {
+    0x0, 0x80000000, 0x0, 0xffffffff, 0xff800000, 0x7f800000, 0x7f800000, 0x7f800000, 0x7f008000, 0x7f800000,
+    0xfe800000, 0x7f000000, 0x7f380000, 0x7f400000, 0x5f900000, 0xd3900000, 0xff008000, 0x200000, 0xfeb80000, 0x9f800000,
+    0x0, 0x80000000, 0x0, 0xffffffff, 0xff800000, 0x7f800000, 0x7f7fffff, 0x7f7fffff, 0x7f008000, 0x7f7fffff,
+    0xfe800000, 0x7f000000, 0x7f380000, 0x7f3fffff, 0x5f8fffff, 0xd38fffff, 0xff008000, 0x200000, 0xfeb80000, 0x9f800000,
+    0x0, 0x80000000, 0x0, 0xffffffff, 0xff800000, 0x7f800000, 0x7f7fffff, 0x7f7fffff, 0x7f008000, 0x7f7fffff,
+    0xfe800000, 0x7f000000, 0x7f380000, 0x7f3fffff, 0x5f8fffff, 0xd3900000, 0xff008000, 0x200000, 0xfeb80000, 0x9f800000,
+    0x0, 0x80000000, 0x0, 0xffffffff, 0xff800000, 0x7f800000, 0x7f800000, 0x7f800000, 0x7f008000, 0x7f800000,
+    0xfe800000, 0x7f000000, 0x7f380000, 0x7f400000, 0x5f900000, 0xd3900000, 0xff008000, 0x200000, 0xfeb80000, 0x9f800000,
+    0x0, 0x80000000, 0x1, 0xffffffff, 0xff800000, 0x7f800000, 0x7f800000, 0x7f800000, 0x7f008000, 0x7f800000,
+    0xfe800000, 0x7f000001, 0x7f380000, 0x7f400000, 0x5f900000, 0xd38fffff, 0xff008000, 0x200000, 0xfeb80000, 0x9f800000
+};
+
+long *input_data;
+const int count_input = 20;
+const int max_function = 5;
+
+static void builtin_convert_double2float(void)
+{
+  // Setup kernel and buffers
+  int k, i, index_cur;
+  int gpu_data[max_function * count_input] = {0};
+  float diff;
+  char log[256] = {0};
+
+  OCL_CREATE_KERNEL("builtin_convert_double2float");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, count_input * max_function * 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(int), &count_input);
+
+  globals[0] = 1;
+  locals[0] = 1;
+  input_data = new long [count_input] ;
+
+  for(int i =0; i < 6; i++)
+    input_data[i] = doubleHX[i];
+  for(int i = 0; i < 14; i++)
+    input_data[i+6] = *((long *)&doubleX +i);
+
+  clEnqueueWriteBuffer( queue, buf[0], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+
+   // Run the kernel
+  OCL_NDRANGE( 1 );
+
+    clEnqueueReadBuffer( queue, buf[1], CL_TRUE, 0, sizeof(int) * max_function* count_input, gpu_data, 0, NULL, NULL);
+    for (k = 0; (uint)k < count_input*max_function; k++)
+    {
+        OCL_ASSERT(gpu_data[k] == expectResult[k]);
+        if(gpu_data[k] != expectResult[k])
+        {
+            printf("failed at function:%s, index:%d  expect value: %d, but get :%d \n", testFunc[k/count_input], k%count_input, expectResult[k], gpu_data[k]);
+        }
+    }
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_convert_double2float)
+}
-- 
2.7.4