utests: add utest to double version of log, log2, log10

Submitted by rander on March 24, 2017, 1:24 a.m.

Details

Message ID 1490318661-10439-1-git-send-email-rander.wang@intel.com
State New
Headers show
Series "utests: add utest to double version of log, log2, log10" ( rev: 1 ) in Beignet

Not browsing as part of any series.

Commit Message

rander March 24, 2017, 1:24 a.m.
Signed-off-by: rander <rander.wang@intel.com>
---
 kernels/builtin_double_logx.cl | 20 ++++++++++
 utests/CMakeLists.txt          |  3 +-
 utests/builtin_double_logx.cpp | 86 ++++++++++++++++++++++++++++++++++++++++++
 3 files changed, 108 insertions(+), 1 deletion(-)
 create mode 100644 kernels/builtin_double_logx.cl
 create mode 100644 utests/builtin_double_logx.cpp

Patch hide | download patch | download mbox

diff --git a/kernels/builtin_double_logx.cl b/kernels/builtin_double_logx.cl
new file mode 100644
index 0000000..ed8e69e
--- /dev/null
+++ b/kernels/builtin_double_logx.cl
@@ -0,0 +1,20 @@ 
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_double_logx(__global double *X,
+												__global double *Z,
+												int max_input)
+{
+	int i = get_global_id(0);
+	int j;
+	double dfloor;
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = log(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = log2(X[j]);
+
+	for(j = 0; j < max_input; j++)
+		Z[i++] = log10(X[j]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 10b0172..41958e8 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -308,7 +308,8 @@  set (utests_sources
   builtin_convert_int16toDouble.cpp
   builtin_convert_double2float.cpp
   builtin_convert_ulong2double.cpp
-  builtin_double_fabs_floor_ceil_fract.cpp)
+  builtin_double_fabs_floor_ceil_fract.cpp
+  builtin_double_logx.cpp)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/builtin_double_logx.cpp b/utests/builtin_double_logx.cpp
new file mode 100644
index 0000000..51f5412
--- /dev/null
+++ b/utests/builtin_double_logx.cpp
@@ -0,0 +1,86 @@ 
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+namespace{
+
+double argX[] {
+    0.0,
+    -0.0,
+    0x0.0000001p-256,
+    -0x0.0000000001p-1022,
+    0x1.0000001p1,
+    -0x1.ffffffp2,
+    0x1.01fffffffffp32,
+    -0x.10ffffffffp48,
+    0x1.7ffffffp64,
+    -0x.1cp96,
+    0x1.00000fp128
+    -0x1.000000000fp256,
+    0x1.00000000000001p512,
+    -0x1.00000000000001p768,
+    0x1.0000000000000fp1023,
+    -0x1.00000000000001p1023
+};
+
+const char*  testFunc[] =
+{
+    "OVERLOADABLE double log(double x)"
+    "OVERLOADABLE double log2(double x)"
+    "OVERLOADABLE double log10(double x)"
+};
+
+unsigned long expectResult[] {
+    0xfff0000000000000, 0xfff0000000000000, 0xc0689b5252dd9845, 0xfff8000000000000, 0x3fe62e4300fa39ef, 0xfff8000000000000, 0x4036304101a0eae6, 0xfff8000000000000, 
+    0x4046622946ce63ca, 0xfff8000000000000, 0xfff8000000000000, 0x40762e42fefa39ef, 0xfff8000000000000, 0x408628b76e3a7b61, 0xfff8000000000000, 0xfff0000000000000, 
+    0xfff0000000000000, 0xfff0000000000000, 0xc071c00000000000, 0xfff8000000000000, 0x3ff0000001715476, 0xfff8000000000000, 0x4040016fe50b6ee5, 0xfff8000000000000, 
+    0x40502570068aa614, 0xfff8000000000000, 0xfff8000000000000, 0x4080000000000000, 0xfff8000000000000, 0x408ff80000000000, 0xfff8000000000000, 0xfff0000000000000, 
+    0xfff0000000000000, 0xfff0000000000000, 0xc0555f856d70eb57, 0xfff8000000000000, 0x3fd34413525c31b0, 0xfff8000000000000, 0x402345ce4de69c10, 0xfff8000000000000, 
+    0x40337127a1b19985, 0xfff8000000000000, 0xfff8000000000000, 0x40634413509f79ff, 0xfff8000000000000, 0x40733f424bcb5220, 0xfff8000000000000, 0xfff0000000000000
+};
+
+double *input_data = argX;
+const int count_input = 16;
+const int max_function = 3;
+
+static void builtin_double_logx(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_double_logx");
+
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, count_input * sizeof(double), NULL);
+  OCL_CREATE_BUFFER(buf[1], 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(int), &count_input);
+
+  globals[0] = 1;
+  locals[0] = 1;
+
+
+  clEnqueueWriteBuffer( queue, buf[0], CL_TRUE, 0, count_input * sizeof(double), input_data, 0, NULL, NULL);
+
+   // Run the kernel
+  OCL_NDRANGE( 1 );
+
+    clEnqueueReadBuffer( queue, buf[1], CL_TRUE, 0, sizeof(double) * max_function * count_input, gpu_data, 0, NULL, NULL);
+
+    int index = 0;
+    for (k = 0; (uint)k < count_input*max_function; k++)
+    {
+       OCL_ASSERT(abs(gpu_data[k] - expectResult[k]) < 2);
+        if(abs(gpu_data[k] - expectResult[k]) > 2)
+        {
+            printf("failed at function:%s, index:%d  expect value: %lx, but get :%lx \n", testFunc[k/count_input], k%count_input, expectResult[k], gpu_data[k]);
+        }
+     }
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_double_logx)
+}

Comments

> +static void builtin_double_logx(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_double_logx");

> +


I just thought of a problem. Seems that you need to check whether cl_khr_fp64 extension is supported or not.
For older generation hardware. It is not possible to support double. Then this utest will certainly fail.
So, I think you need to add the this extension check for all the double test cases.

Thanks!
Ruiling
As I know ivyBrigde support double, it there anything wrong?

-----Original Message-----
From: Song, Ruiling 

Sent: Wednesday, March 29, 2017 11:06 AM
To: Wang, Rander <rander.wang@intel.com>; beignet@freedesktop.org
Cc: Wang, Rander <rander.wang@intel.com>
Subject: RE: [Beignet] [PATCH] utests: add utest to double version of log, log2, log10

> +static void builtin_double_logx(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_double_logx");

> +


I just thought of a problem. Seems that you need to check whether cl_khr_fp64 extension is supported or not.
For older generation hardware. It is not possible to support double. Then this utest will certainly fail.
So, I think you need to add the this extension check for all the double test cases.

Thanks!
Ruiling
> -----Original Message-----

> From: Wang, Rander

> Sent: Wednesday, March 29, 2017 1:12 PM

> To: Song, Ruiling <ruiling.song@intel.com>; beignet@freedesktop.org

> Subject: RE: [Beignet] [PATCH] utests: add utest to double version of log, log2,

> log10

> 

> As I know ivyBrigde support double, it there anything wrong?

IVB did support many double operations. But division operation on double is not supported until GEN8.
And what's more, the double add/mul on IVB and HSW has very low precision which is far from required if I remember correctly.
So, we can only support double from GEN8.

Thanks!
Ruiling
> 

> -----Original Message-----

> From: Song, Ruiling

> Sent: Wednesday, March 29, 2017 11:06 AM

> To: Wang, Rander <rander.wang@intel.com>; beignet@freedesktop.org

> Cc: Wang, Rander <rander.wang@intel.com>

> Subject: RE: [Beignet] [PATCH] utests: add utest to double version of log, log2,

> log10

> 

> > +static void builtin_double_logx(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_double_logx");

> > +

> 

> I just thought of a problem. Seems that you need to check whether cl_khr_fp64

> extension is supported or not.

> For older generation hardware. It is not possible to support double. Then this

> utest will certainly fail.

> So, I think you need to add the this extension check for all the double test cases.

> 

> Thanks!

> Ruiling