[Beignet] [PATCH 2/2] test if register allocation and 64-bit reading are fixed

[Homer Hsing]

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

diff --git a/kernels/compiler_double_4.cl b/kernels/compiler_double_4.cl
new file mode 100644
index 0000000..e5e46f9
--- /dev/null
+++ b/kernels/compiler_double_4.cl
@@ -0,0 +1,5 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+kernel void compiler_double_4(global double *src1, global double *src2, global double *dst) {
+  int i = get_global_id(0);
+  dst[i] = src1[i] + src2[i];
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 553e654..ef924cd 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -117,6 +117,7 @@ set (utests_sources
   compiler_double.cpp
   compiler_double_2.cpp
   compiler_double_3.cpp
+  compiler_double_4.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp
diff --git a/utests/compiler_double_4.cpp b/utests/compiler_double_4.cpp
new file mode 100644
index 0000000..cb25bd4
--- /dev/null
+++ b/utests/compiler_double_4.cpp
@@ -0,0 +1,40 @@
+#include <cmath>
+#include "utest_helper.hpp"
+
+void compiler_double_4(void)
+{
+  const size_t n = 16;
+  double cpu_src1[n], cpu_src2[n];
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_double_4");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(double), 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;
+
+  // Run random tests
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    cpu_src1[i] = ((double*)buf_data[0])[i] = rand() * 1e-2;
+    cpu_src2[i] = ((double*)buf_data[1])[i] = rand() * 1e-2;
+  }
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(2);
+  for (int32_t i = 0; i < (int32_t) n; ++i)
+    OCL_ASSERT(fabs(((double*)buf_data[2])[i] - cpu_src1[i] - cpu_src2[i]) < 1e-4);
+  OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_double_4);
-- 
1.8.1.2