[Beignet] [PATCH 1/2] padding enough temporary registers in 64-bit IO

[Xing, Homer]

Ping for review.

http://lists.freedesktop.org/archives/beignet/2013-July/001007.html

-----Original Message-----
From: beignet-bounces+homer.xing=intel.com at lists.freedesktop.org [mailto:beignet-bounces+homer.xing=intel.com at lists.freedesktop.org] On Behalf Of Homer Hsing
Sent: Wednesday, July 24, 2013 10:39 AM
To: beignet at lists.freedesktop.org
Subject: [Beignet] [PATCH 2/2] test previous patch


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 95f3d9b..b2670bf 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -112,6 +112,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

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet