[Beignet] [PATCH 2/2] utests: add a double precision check test case.

[Zhigang Gong]

Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
---
 kernels/double_precision_check.cl    | 11 +++++++++
 utests/CMakeLists.txt                |  1 +
 utests/compiler_double_precision.cpp | 43 ++++++++++++++++++++++++++++++++++++
 3 files changed, 55 insertions(+)
 create mode 100644 kernels/double_precision_check.cl
 create mode 100644 utests/compiler_double_precision.cpp

diff --git a/kernels/double_precision_check.cl b/kernels/double_precision_check.cl
new file mode 100644
index 0000000..6dd1123
--- /dev/null
+++ b/kernels/double_precision_check.cl
@@ -0,0 +1,11 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void
+double_precision_check(__global float* src, __global float* dst)
+{
+  int id = (int)get_global_id(0);
+  double d0 = 0.12345678912345678 + src[1];
+  double d1 = 0.12355678912345688 + src[0];
+  float rem = d1 - d0;
+  dst[id] = rem;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 9ad08c9..1c523cb 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -167,6 +167,7 @@ set (utests_sources
   compiler_getelementptr_bitcast.cpp
   compiler_simd_any.cpp
   compiler_simd_all.cpp
+  compiler_double_precision.cpp
   load_program_from_bin.cpp
   profiling_exec.cpp
   enqueue_copy_buf.cpp
diff --git a/utests/compiler_double_precision.cpp b/utests/compiler_double_precision.cpp
new file mode 100644
index 0000000..d422132
--- /dev/null
+++ b/utests/compiler_double_precision.cpp
@@ -0,0 +1,43 @@
+#include "utest_helper.hpp"
+#include <math.h>
+
+static void double_precision_check(void)
+{
+  const size_t n = 16; //8192 * 4;
+
+  double d0 = 0.12345678912345678;
+  double d1 = 0.12355678922345678;
+  float cpu_result = d1 - d0;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("double_precision_check");
+  //OCL_CREATE_KERNEL("compiler_array");
+  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((float*)buf_data[0])[i] = 0;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+  free(buf_data[0]);
+  buf_data[0] = NULL;
+
+  // Run the kernel
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  bool precisionOK = true;
+  for (uint32_t i = 0; i < n; ++i) {
+    float error = ((float*)buf_data[1])[i] - cpu_result;
+    if (error - cpu_result != 0)
+      precisionOK = false;
+    OCL_ASSERT((fabs(error) < 1e-4));
+  }
+  if (!precisionOK)
+    printf("\n  - WARN: GPU doesn't have correct double precision. Got %.7G, expected %.7G\n", ((float*)buf_data[1])[0], cpu_result);
+}
+
+MAKE_UTEST_FROM_FUNCTION(double_precision_check);
-- 
1.8.3.2