[Beignet] [PATCH] utests: add utest to convert_double_rd(ulong)

[rander]

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

diff --git a/kernels/builtin_convert_ulong2double.cl b/kernels/builtin_convert_ulong2double.cl
new file mode 100644
index 0000000..6c78e1b
--- /dev/null
+++ b/kernels/builtin_convert_ulong2double.cl
@@ -0,0 +1,25 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_convert_ulong2double(__global ulong *X,
+												__global double *Z,
+												int max_input)
+{
+	int i = get_global_id(0);
+	int j;
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_double(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_double_rtz(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_double_rtn(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_double_rte(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = convert_double_rtp(X[j]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index d76bba7..0f69a0e 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -306,7 +306,8 @@ set (utests_sources
   builtin_convert_double2int64.cpp
   builtin_convert_int8toDouble.cpp
   builtin_convert_int16toDouble.cpp
-  builtin_convert_double2float.cpp)
+  builtin_convert_double2float.cpp
+  builtin_convert_ulong2double.cpp)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/builtin_convert_ulong2double.cpp b/utests/builtin_convert_ulong2double.cpp
new file mode 100644
index 0000000..4706541
--- /dev/null
+++ b/utests/builtin_convert_ulong2double.cpp
@@ -0,0 +1,91 @@
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+namespace{
+    unsigned long ulong[] = {
+        0x0,
+        0x8000000000000000,
+        0x0000000000000001,
+        0xFFFFFFFFFFFFFFFF,
+        0xFFF0000000000000,
+        0x7FF0000000000000,
+        0x00000000FFFFFFFF,
+        0x000FFFFFFFFFFFFF,
+        0x0010000000000000,
+        0x0018000000000000,
+        0x001F000000000000,
+        0x0017000000000000,
+        0x0000000FFFFFFFFF,
+        0x000000FFFFFFFFFF,
+        0x0000000000000002,
+        0x0000000000000010,
+};
+
+
+const char*  testFunc[] =
+{
+    " char convert_double(ulong x)",
+    " char convert_double_sat(ulong x)",
+    " char convert_double_rtz(ulong x)",
+    " char convert_double_rtn(ulong x)",
+    " char convert_double_rte(ulong x)",
+};
+
+unsigned long expectResult[] = {
+    0x0, 0x43e0000000000000, 0x3ff0000000000000, 0x43f0000000000000, 0x43effe0000000000, 0x43dffc0000000000, 0x41efffffffe00000, 0x432ffffffffffffe,
+    0x4330000000000000,0x4338000000000000, 0x433f000000000000, 0x4337000000000000, 0x422ffffffffe0000, 0x426fffffffffe000, 0x4000000000000000, 0x4030000000000000,
+    0x0,0x43e0000000000000, 0x3ff0000000000000, 0x43efffffffffffff, 0x43effe0000000000, 0x43dffc0000000000, 0x41efffffffe00000, 0x432ffffffffffffe,
+    0x4330000000000000,0x4338000000000000, 0x433f000000000000, 0x4337000000000000, 0x422ffffffffe0000, 0x426fffffffffe000, 0x4000000000000000, 0x4030000000000000,
+    0x0,0x43e0000000000000, 0x3ff0000000000000, 0x43efffffffffffff, 0x43effe0000000000, 0x43dffc0000000000, 0x41efffffffe00000, 0x432ffffffffffffe,
+    0x4330000000000000,0x4338000000000000, 0x433f000000000000, 0x4337000000000000, 0x422ffffffffe0000, 0x426fffffffffe000, 0x4000000000000000, 0x4030000000000000,
+    0x0,0x43e0000000000000, 0x3ff0000000000000, 0x43f0000000000000, 0x43effe0000000000, 0x43dffc0000000000, 0x41efffffffe00000, 0x432ffffffffffffe,
+    0x4330000000000000,0x4338000000000000, 0x433f000000000000, 0x4337000000000000, 0x422ffffffffe0000, 0x426fffffffffe000, 0x4000000000000000, 0x4030000000000000,
+    0x0,0x43e0000000000000, 0x3ff0000000000000, 0x43f0000000000000, 0x43effe0000000000, 0x43dffc0000000000, 0x41efffffffe00000, 0x432ffffffffffffe,
+    0x4330000000000000,0x4338000000000000, 0x433f000000000000, 0x4337000000000000, 0x422ffffffffe0000, 0x426fffffffffe000, 0x4000000000000000, 0x4030000000000000
+
+};
+
+unsigned long *input_data;
+const int count_input = 16;
+const int max_function = 5;
+
+static void builtin_convert_ulong2double(void)
+{
+  // Setup kernel and buffers
+  int k, i, index_cur;
+  unsigned long gpu_data[max_function * count_input] = {0};
+  float diff;
+  char log[256] = {0};
+
+  OCL_CREATE_KERNEL("builtin_convert_ulong2double");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, count_input * max_function * sizeof(long), NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, count_input * max_function * sizeof(long), 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 = ulong;
+
+  clEnqueueWriteBuffer( queue, buf[0], CL_TRUE, 0, count_input * max_function * sizeof(long), input_data, 0, NULL, NULL);
+
+   // Run the kernel
+  OCL_NDRANGE( 1 );
+
+    clEnqueueReadBuffer( queue, buf[1], CL_TRUE, 0, sizeof(long) * 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: %lx, but get :%lx \n", testFunc[k/count_input], k%count_input, expectResult[k], gpu_data[k]);
+        }
+    }
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_convert_ulong2double)
+}
-- 
2.7.4