test if register allocation and 64-bit reading are fixed

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>

diff --git a/kernels/compiler_double_4.cl b/kernels/compiler_double_4.cl
new file mode 100644 (file)
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 4cdb449..ceb5e08 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -120,6 +120,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 (file)
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);