[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