[Beignet] [PATCH] utest: add utest to common function of double

rander rander.wang at intel.com
Mon Mar 13 06:27:27 UTC 2017


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

diff --git a/kernels/builtin_commonFunc_fp64.cl b/kernels/builtin_commonFunc_fp64.cl
new file mode 100644
index 0000000..6128de3
--- /dev/null
+++ b/kernels/builtin_commonFunc_fp64.cl
@@ -0,0 +1,41 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_commonFunc_fp64(__global double *X,
+									__global double *Y,
+									__global double *XPlus,
+									__global double *XSmooth,
+									__global double *XSign,
+									__global double *Z,
+									int max_input)
+{
+	int i = get_global_id(0);
+	int j;
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = step(X[j], Y[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = max(X[j], Y[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = min(X[j], Y[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = mix(X[j], Y[j], 1/max_input);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = clamp(Y[j], X[j], XPlus[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = degrees(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = radians(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = smoothstep(X[j], XPlus[j], XSmooth[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = sign(XSign[j]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 6e731af..a4ab0ab 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -298,7 +298,8 @@ set (utests_sources
   runtime_pipe_query.cpp
   compiler_pipe_builtin.cpp
   compiler_device_enqueue.cpp
-  builtin_relation_fp64)
+  builtin_relation_fp64.cpp
+  builtin_commonFunc_fp64.cpp)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/builtin_commonFunc_fp64.cpp b/utests/builtin_commonFunc_fp64.cpp
new file mode 100644
index 0000000..a590ca9
--- /dev/null
+++ b/utests/builtin_commonFunc_fp64.cpp
@@ -0,0 +1,226 @@
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+#define udebug 0
+
+#define FLT_MAX 0x1.fffffep127f
+#define FLT_MIN ldexpf(1.0,-126)
+#define FLT_ULP  (1.0e-6f)
+
+#define printf_c(...) \
+{\
+  printf("\033[1m\033[40;31m");\
+  printf( __VA_ARGS__ );\
+  printf("\033[0m");\
+}
+
+namespace{
+double doubleX[] = {
+        0.0,
+        0x1.0p0,
+        -0x1.1p2,
+        0x1.8p5,
+        -0x1.FEp10,
+        0x1.FFFFFFFFFFp12,
+        -0x1.FFFFFFFFFFp12,
+        0x1.00000001p22,
+        -0x1.00000001p22,
+        0x1.7777777777p33,
+        -0x1.7777777777p33,
+        0x0.00000000001p44,
+        -0x0.00000000001p44,
+        0x1.FFFFFFFFFFp128,
+        -0x1.FFFFFFFFFFp128,
+        0x1.FFFFFFFFFFp1023
+};
+
+double doubleY[] = {
+        -0x1.00000001p22,
+        0.0,
+        -0x1.1p2,
+        0x1.FFFFFFFFFFp128,
+        0x1.8p5,
+        -0x1.FEp10,
+        0x1.FFFFFFFFFFp12,
+        0x1.00000001p22,
+        0x1.FFFFFFFFFFp1023,
+        0x1.7777777777p33,
+        0x1.0p0,
+        -0x1.7777777777p33,
+        -0x1.FFFFFFFFFFp12,
+        0x0.00000000001p44,
+        -0x0.00000000001p44,
+        -0x1.FFFFFFFFFFp128,
+};
+
+double doubleXPlus[] = {
+        0.0 + 0.1,
+        0x1.0p0*2,
+        -0x1.1p2 + 3,
+        0x1.8p5 + 5,
+        -0x1.FEp10 + 10,
+        0x1.FFFFFFFFFFp12 + 12,
+        -0x1.FFFFFFFFFFp12 + 12,
+        0x1.00000001p22 + 22,
+        -0x1.00000001p22 + 22,
+        0x1.7777777777p33 + 33,
+        -0x1.7777777777p33 + 33,
+        0x0.00000000001p44 + 44,
+        -0x0.00000000001p44 + 44,
+        0x1.FFFFFFFFFFp128 + 128,
+        -0x1.FFFFFFFFFFp128 + 128,
+        0x1.FFFFFFFFFFp1023 + 2048
+};
+
+double doubleXSmooth[] = {
+        1.0,
+        0x1.0p0 - 1,
+        -0x1.1p2 +1,
+        0x1.8p5 + 2,
+        -0x1.FEp10 + 5,
+        0x1.FFFFFFFFFFp12 + 6,
+        -0x1.FFFFFFFFFFp12 + 5,
+        0x1.00000001p22 + 10,
+        -0x1.00000001p22 + 11,
+        0x1.7777777777p33 + 20,
+        -0x1.7777777777p33 + 23,
+        0x0.00000000001p44 + 21,
+        -0x0.00000000001p44 + 25,
+        0x1.FFFFFFFFFFp128 + 60,
+        -0x1.FFFFFFFFFFp128 + 61,
+        0x1.FFFFFFFFFFp1023 + 1000
+};
+
+unsigned long doubleBinaryX[] = {
+        0x0UL, //+0
+        0x8000000000000000UL, //-0
+        0x7FFFFFFFFFFFFFFFUL, //NAN
+        0x0000000000000001UL,  // DENORM
+        0x8000000000000001UL,  //-DENORM
+        0x7FE0000000000001UL,  // normal
+        0x7FE0000000000002UL,  // normal
+        0xFFF0000000000000UL,  //-INFINITY
+        0xFFE0000000000001UL,  // normal
+        0x3F80000000000000UL,   // normal
+        0x7FF0000000000000UL, //+INFINITY
+        0x5671230000000000UL,  // normal
+        0x1234560000000000UL,  // normal
+        0xA120234000000000UL, // normal
+        0xC451200088771111UL, // normal
+        0x1000241234110000UL // normal
+};
+
+const char*  testFunc[] =
+{
+    " double step(double edge, double x)",
+    " double max(double a, double b)",
+    " double min(double a, double b)",
+    " double mix(double x, double y, double a)",
+    " double clamp(double v, double l, double u)",
+    " double smoothstep(double e0, double e1, double x)",
+    " double degrees(double radians)",
+    " double radians(double degrees)",
+    " double sign(double x)"
+};
+
+unsigned long expectResult[] = {
+    0x0               ,    0x0               ,    0x3ff0000000000000,    0x3ff0000000000000,
+    0x3ff0000000000000,    0x0               ,    0x3ff0000000000000,    0x3ff0000000000000,
+    0x3ff0000000000000,    0x3ff0000000000000,    0x3ff0000000000000,    0x0               ,
+    0x0               ,    0x0               ,    0x3ff0000000000000,    0x0               ,
+    0x0               ,    0x3ff0000000000000,    0xc011000000000000,    0x47fffffffffff000,
+    0x4048000000000000,    0x40bffffffffff000,    0x40bffffffffff000,    0x4150000000100000,
+    0x7feffffffffff000,    0x4207777777777000,    0x3ff0000000000000,    0x3ff0000000000000,
+    0xbff0000000000000,    0x47fffffffffff000,    0xbff0000000000000,    0x7feffffffffff000,
+    0xc150000000100000,    0x0               ,    0xc011000000000000,    0x4048000000000000,
+    0xc09fe00000000000,    0xc09fe00000000000,    0xc0bffffffffff000,    0x4150000000100000,
+    0xc150000000100000,    0x4207777777777000,    0xc207777777777000,    0xc207777777777000,
+    0xc0bffffffffff000,    0x3ff0000000000000,    0xc7fffffffffff000,    0xc7fffffffffff000,
+    0x0               ,    0x3ff0000000000000,    0xc011000000000000,    0x4048000000000000,
+    0xc09fe00000000000,    0x40bffffffffff000,    0xc0bffffffffff000,    0x4150000000100000,
+    0xc150000000100000,    0x4207777777777000,    0xc207777777777000,    0x3ff0000000000000,
+    0xbff0000000000000,    0x47fffffffffff000,    0xc7fffffffffff000,    0x7feffffffffff000,
+    0x0               ,    0x3ff0000000000000,    0xc011000000000000,    0x404a800000000000,
+    0xc09fb80000000000,    0x40bffffffffff000,    0xc0bff3fffffff000,    0x4150000000100000,
+    0xc14ffff500200000,    0x4207777777777000,    0xc2077777766f7000,    0x3ff0000000000000,
+    0xbff0000000000000,    0x47fffffffffff000,    0xc7fffffffffff000,    0x7feffffffffff000,
+    0x0               ,    0x404ca5dc1a63c284,    0xc06e7039dc09feac,    0x40a57c6513cad1e3,
+    0xc0fc89363e495ec1,    0x411ca5dc1a63b431,    0xc11ca5dc1a63b431,    0x41aca5dc1a806860,
+    0xc1aca5dc1a806860,    0x42650229f138107e,    0xc2650229f138107e,    0x404ca5dc1a63c284,
+    0xc04ca5dc1a63c284,    0x485ca5dc1a63b431,    0xc85ca5dc1a63b431,    0x7ff0000000000000,
+    0x0               ,    0x3f91df46a2529ce1,    0xbfb2fd3b0c77c6af,    0x3feacee9f37beb52,
+    0xc041cd675bb04a44,    0x4061df46a25293f1,    0xc061df46a25293f1,    0x40f1df46a2647c28,
+    0xc0f1df46a2647c28,    0x41aa366798bd6648,    0xc1aa366798bd6648,    0x3f91df46a2529ce1,
+    0xbf91df46a2529ce1,    0x47a1df46a25293f1,    0xc7a1df46a25293f1,    0x7f91df46a25293f1,
+    0x3ff0000000000000,    0x0               ,    0x3fd097b425ed097b,    0x3fd6872b020c49bc,
+    0x3fe0000000000000,    0x3fe0000000000000,    0x3fd812f684bda130,    0x3fdba5fc89b68704,
+    0x3fe0000000000000,    0x3fe503b9c2f370ee,    0x3fe8f728de32d91d,    0x3fddd1d6d708453f,
+    0x3fe340a006279f11,    0x3ff0000000000000,    0x3ff0000000000000,    0x3ff0000000000000,
+    0x0               ,    0x8000000000000000,    0x0               ,    0x3ff0000000000000,
+    0xbff0000000000000,    0x3ff0000000000000,    0x3ff0000000000000,    0xbff0000000000000,
+    0xbff0000000000000,    0x3ff0000000000000,    0x3ff0000000000000,    0x3ff0000000000000,
+    0x3ff0000000000000,    0xbff0000000000000,    0xbff0000000000000,    0x3ff0000000000000
+};
+
+double *input_data;
+const int count_input = 16;
+const int max_function = 9;
+
+static void builtin_commonFunc_fp64(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_commonFunc_fp64");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[2], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[3], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[4], CL_MEM_READ_WRITE, count_input * max_function * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[5], CL_MEM_READ_WRITE, count_input * max_function * 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]);
+  OCL_SET_ARG(3, sizeof(cl_mem), &buf[3]);
+  OCL_SET_ARG(4, sizeof(cl_mem), &buf[4]);
+  OCL_SET_ARG(5, sizeof(cl_mem), &buf[5]);
+  OCL_SET_ARG(6, sizeof(int), &count_input);
+
+  globals[0] = 1;
+  locals[0] = 1;
+
+  input_data = (double *)doubleX;
+  clEnqueueWriteBuffer( queue, buf[0], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+  input_data = (double *)doubleY;
+  clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+  input_data = (double *)doubleXPlus;
+  clEnqueueWriteBuffer( queue, buf[2], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+  input_data = (double *)doubleXSmooth;
+  clEnqueueWriteBuffer( queue, buf[3], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+  input_data = (double *)doubleBinaryX;
+  clEnqueueWriteBuffer( queue, buf[4], CL_TRUE, 0, count_input * max_function * sizeof(double), input_data, 0, NULL, NULL);
+
+   // Run the kernel
+  OCL_NDRANGE( 1 );
+
+  clEnqueueReadBuffer( queue, buf[5], CL_TRUE, 0, sizeof(long) * max_function * count_input, gpu_data, 0, NULL, NULL);
+
+  printf("\n");
+
+  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,  expect value: %lx, but get :%lx \n", testFunc[k/count_input], expectResult[k], gpu_data[k]);
+        }
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_commonFunc_fp64)
+}
-- 
2.7.4



More information about the Beignet mailing list