[Beignet] [PATCH] utests: add utest to double type of simple math functions
rander
rander.wang at intel.com
Mon Mar 27 08:38:50 UTC 2017
double copysign(double x, double y)
double fdim(double x, double y)
double fmax(double x, double y)
double fmin(double x, double y)
double frexp(double x)
double mad(double x, double y, , double z)
double maxmag(double x, double y)
double minmag(double x, double y)
double nan(double x)
double nextafter(double x)
double rint(double x)
double round(double x)
double trunc(double x)
Signed-off-by: rander <rander.wang at intel.com>
---
kernels/builtin_double_simpleMathFuncs.cl | 52 +++++++++++++
utests/CMakeLists.txt | 3 +-
utests/builtin_double_simpleMathFuncs.cpp | 124 ++++++++++++++++++++++++++++++
3 files changed, 178 insertions(+), 1 deletion(-)
create mode 100644 kernels/builtin_double_simpleMathFuncs.cl
create mode 100644 utests/builtin_double_simpleMathFuncs.cpp
diff --git a/kernels/builtin_double_simpleMathFuncs.cl b/kernels/builtin_double_simpleMathFuncs.cl
new file mode 100644
index 0000000..d4ffebf
--- /dev/null
+++ b/kernels/builtin_double_simpleMathFuncs.cl
@@ -0,0 +1,52 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_double_simpleMathFuncs( __global double *X,
+ __global double *Y,
+ __global double *Z,
+ int max_input)
+{
+ int i = get_global_id(0);
+ int j;
+ int expVal;
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = copysign(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = fdim(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = fmax(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = fmin(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = frexp(X[j], &expVal);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = mad(X[j], Y[j], X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = maxmag(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = minmag(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = nan(as_ulong(X[j]));
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = nextafter(X[j], Y[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = rint(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = round(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = trunc(X[j]);
+
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 41958e8..a6506ff 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -309,7 +309,8 @@ set (utests_sources
builtin_convert_double2float.cpp
builtin_convert_ulong2double.cpp
builtin_double_fabs_floor_ceil_fract.cpp
- builtin_double_logx.cpp)
+ builtin_double_logx.cpp
+ builtin_double_simpleMathFuncs.cpp)
if (LLVM_VERSION_NODOT VERSION_GREATER 34)
SET(utests_sources
diff --git a/utests/builtin_double_simpleMathFuncs.cpp b/utests/builtin_double_simpleMathFuncs.cpp
new file mode 100644
index 0000000..0681e82
--- /dev/null
+++ b/utests/builtin_double_simpleMathFuncs.cpp
@@ -0,0 +1,124 @@
+#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.01fffffffffp12,
+ -0x.10ffffffffp14,
+ 0x1.7ffffffp8,
+ -0x.1cp10,
+ 0x1.00000fp9
+ -0x1.000000000fp10,
+};
+
+double argY[] {
+ -0x1.000000000fp10,
+ 0x1.00000fp9
+ -0x.1cp10,
+ 0x1.7ffffffp8,
+ -0x.10ffffffffp14,
+ -0x1.ffffffp2,
+ 0x1.0000001p1,
+ -0x0.0000000001p-1022,
+ 0x0.0000001p-256,
+ -0.0,
+ 0.0
+};
+
+const char* testFunc[] =
+{
+ "OVERLOADABLE double copysign(double x, double y)"
+ "OVERLOADABLE double fdim(double x, double y)"
+ "OVERLOADABLE double fmax(double x, double y)"
+ "OVERLOADABLE double fmin(double x, double y)"
+ "OVERLOADABLE double frexp(double x)"
+ "OVERLOADABLE double mad(double x, double y, , double z)"
+ "OVERLOADABLE double maxmag(double x, double y)"
+ "OVERLOADABLE double minmag(double x, double y)"
+ "OVERLOADABLE double nan(double x)"
+ "OVERLOADABLE double nextafter(double x)"
+ "OVERLOADABLE double rint(double x)"
+ "OVERLOADABLE double round(double x)"
+ "OVERLOADABLE double trunc(double x)"
+};
+
+unsigned long expectResult[] {
+ 0x8000000000000000, 0x0, 0x2e30000000000000, 0x8000000000001000, 0xc000000001000000, 0x401ffffff0000000, 0xc0b01fffffffff00, 0x4090ffffffff0000,
+ 0xc077ffffff000000, 0x405c000000000000, 0x407ffffe2003c000, 0x0, 0x409000000000f000, 0x0, 0x0, 0x4090ffffffff0000,
+ 0x4023fffff8400000, 0x0, 0x40b01fffffffff00, 0x0, 0x4077ffffff000000, 0x0, 0x0, 0x0,
+ 0x0, 0x40790001e0000000, 0x4077ffffff000000, 0x8000000000001000, 0x4000000001000000, 0x4000000001000000, 0x40b01fffffffff00, 0x2e30000000000000,
+ 0x4077ffffff000000, 0x0, 0x0, 0x0, 0xc09000000000f000, 0x8000000000000000, 0x2e30000000000000, 0xc090ffffffff0000,
+ 0xc01ffffff0000000, 0xc01ffffff0000000, 0x8000000000001000, 0xc090ffffffff0000, 0x8000000000000000, 0xc05c000000000000, 0xc07ffffe2003c000, 0x0,
+ 0x0, 0x8000000000000000, 0x3fe0000000000000, 0xbfe0000000000000, 0x3fe0000001000000, 0xbfeffffff0000000, 0x3fe01fffffffff00, 0xbfe0ffffffff0000,
+ 0x3fe7ffffff000000, 0xbfec000000000000, 0xbfeffffe2003c000, 0x0, 0x0, 0x8000000000000000, 0x2eb80fffff000000, 0x43f000,
+ 0xc02bfffff1bfffff, 0xc037fffff5000000, 0x40b01fffffffff00, 0xc090ffffffff0000, 0x4077ffffff000000, 0xc05c000000000000, 0xc07ffffe2003c000, 0x0,
+ 0xc09000000000f000, 0x40790001e0000000, 0x4077ffffff000000, 0xc090ffffffff0000, 0xc01ffffff0000000, 0xc01ffffff0000000, 0x40b01fffffffff00, 0xc090ffffffff0000,
+ 0x4077ffffff000000, 0xc05c000000000000, 0xc07ffffe2003c000, 0x0, 0x0, 0x8000000000000000, 0x2e30000000000000, 0x8000000000001000,
+ 0x4000000001000000, 0x4000000001000000, 0x8000000000001000, 0x2e30000000000000, 0x8000000000000000, 0x0, 0x0, 0x0,
+ 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001,
+ 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x7ff0000000000001, 0x8000000000000001, 0x1, 0x2e30000000000001, 0x8000000000001001,
+ 0x4000000000ffffff, 0xc01fffffefffffff, 0x40b01ffffffffeff, 0xc090fffffffeffff, 0x4077fffffeffffff, 0xc05bffffffffffff, 0xc07ffffe2003bfff, 0x0,
+ 0x0, 0x8000000000000000, 0x0, 0x8000000000000000, 0x4000000000000000, 0xc020000000000000, 0x40b0200000000000, 0xc091000000000000,
+ 0x4078000000000000, 0xc05c000000000000, 0xc080000000000000, 0x0, 0x0, 0x8000000000000000, 0x0, 0x8000000000000000,
+ 0x4000000000000000, 0xc020000000000000, 0x40b0200000000000, 0xc091000000000000, 0x4078000000000000, 0xc05c000000000000, 0xc080000000000000, 0x0,
+ 0x0, 0x8000000000000000, 0x0, 0x8000000000000000, 0x4000000000000000, 0xc01c000000000000, 0x40b01f0000000000, 0xc090fc0000000000,
+ 0x4077f00000000000, 0xc05c000000000000, 0xc07ff00000000000, 0x0
+};
+
+double *input_data = argX;
+const int count_input = 12;
+const int max_function = 13;
+
+static void builtin_double_simpleMathFuncs(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_simpleMathFuncs");
+
+ 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 * sizeof(double), NULL);
+ OCL_CREATE_BUFFER(buf[2], 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(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);
+ clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, count_input * sizeof(double), argY, 0, NULL, NULL);
+
+ // Run the kernel
+ OCL_NDRANGE( 1 );
+
+ clEnqueueReadBuffer( queue, buf[2], 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((gpu_data[k] - expectResult[k]) == 0);
+ if((gpu_data[k] - expectResult[k]) != 0)
+ {
+ 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_simpleMathFuncs)
+}
--
2.7.4
More information about the Beignet
mailing list