[Beignet] [PATCH] utests: add utests to trigonometric of double version
rander
rander.wang at intel.com
Thu Mar 30 05:21:50 UTC 2017
there are 20 trigonometric functions, test them together
Signed-off-by: rander <rander.wang at intel.com>
---
kernels/builtin_double_trigonometric_function.cl | 71 ++++++++++++
utests/CMakeLists.txt | 3 +-
utests/builtin_double_trigonometric_function.cpp | 139 +++++++++++++++++++++++
3 files changed, 212 insertions(+), 1 deletion(-)
create mode 100644 kernels/builtin_double_trigonometric_function.cl
create mode 100644 utests/builtin_double_trigonometric_function.cpp
diff --git a/kernels/builtin_double_trigonometric_function.cl b/kernels/builtin_double_trigonometric_function.cl
new file mode 100644
index 0000000..a84e4dc
--- /dev/null
+++ b/kernels/builtin_double_trigonometric_function.cl
@@ -0,0 +1,71 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_double_trigonometric_function(__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++] = acos(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = acospi(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = acosh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = asin(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = asinpi(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = asinh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = atan(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = atan2(X[j], 1.0);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = atanh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = atanpi(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = atan2pi(X[j], 1.0);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = cos(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = cosh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = cospi(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = sin(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = sinh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = sinpi(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = tan(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = tanh(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = tanpi(X[j]);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index a6506ff..c42fedc 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -310,7 +310,8 @@ set (utests_sources
builtin_convert_ulong2double.cpp
builtin_double_fabs_floor_ceil_fract.cpp
builtin_double_logx.cpp
- builtin_double_simpleMathFuncs.cpp)
+ builtin_double_simpleMathFuncs.cpp
+ builtin_double_trigonometric_function.cpp)
if (LLVM_VERSION_NODOT VERSION_GREATER 34)
SET(utests_sources
diff --git a/utests/builtin_double_trigonometric_function.cpp b/utests/builtin_double_trigonometric_function.cpp
new file mode 100644
index 0000000..de9b889
--- /dev/null
+++ b/utests/builtin_double_trigonometric_function.cpp
@@ -0,0 +1,139 @@
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+namespace{
+
+double argX[] {
+ 0.0,
+ -0.0,
+ 0x1.0p0,
+ -0x1.0p0,
+ 0x1.0p-1,
+ -0x1.0p-1,
+ 0x1.1p-1,
+ -0x1.1p-1,
+ 0x1.0p-2,
+ -0x1.0p-2,
+ 0x1.1p-2,
+ -0x1.1p-2,
+ 0x1.0p-3,
+ -0x1.0p-3,
+ 0x1.1p-3,
+ -0x1.1p-3
+};
+
+const char* testFunc[] =
+{
+ " double acos(double x)"
+ " double acospi(double x)"
+ " double acosh(double x)"
+ " double asin(double x)"
+ " double asinpi(double x)"
+ " double asinh(double x)"
+ " double atan(double x)"
+ " double atan2(double x, double y)"
+ " double atanh(double x)"
+ " double atanpi(double x)"
+ " double atan2pi(double x, double y)"
+ " double cos(double x)"
+ " double cosh(double x)"
+ " double cospi(double x)"
+ " double sin(double x)"
+ " double sinh(double x)"
+ " double sinpi(double x)"
+ " double tan(double x)"
+ " double tanh(double x)"
+ " double tanpi(double x)"
+};
+
+unsigned long expectResult[] {
+ 0x3ff921fb54442d18, 0x3ff921fb54442d18, 0x0, 0x400921fb54442d18, 0x3ff0c152382d7365, 0x4000c152382d7366, 0x3ff02be9ce0b87cd, 0x40010c066d3e6932,
+ 0x3ff51700e0c14b25, 0x3ffd2cf5c7c70f0b, 0x3ff4d4c369e418c6, 0x3ffd6f333ea4416a, 0x3ff720a392c1d954, 0x3ffb235315c680dc, 0x3ff7005ea0cf9800, 0x3ffb439807b8c230,
+ 0x3fe0000000000000, 0x3fe0000000000000, 0x0, 0x3ff0000000000000, 0x3fd5555555555556, 0x3fe5555555555555, 0x3fd49719f07a5377, 0x3fe5b47307c2d644,
+ 0x3fdada3983e16b22, 0x3fe292e33e0f4a6f, 0x3fda85e299dc401d, 0x3fe2bd0eb311dff1, 0x3fdd726443466d1d, 0x3fe146cdde5cc972, 0x3fdd494e19ceca92, 0x3fe15b58f3189ab7,
+ 0x7ff8000000000000, 0x7ff8000000000000, 0x0, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000,
+ 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000, 0x7ff8000000000000,
+ 0x0, 0x8000000000000000, 0x3ff921fb54442d18, 0xbff921fb54442d18, 0x3fe0c152382d7366, 0xbfe0c152382d7366, 0x3fe1ec230c714a96, 0xbfe1ec230c714a96,
+ 0x3fd02be9ce0b87cd, 0xbfd02be9ce0b87cd, 0x3fd134dfa9805146, 0xbfd134dfa9805146, 0x3fc00abe0c129e1e, 0xbfc00abe0c129e1e, 0x3fc10ce59ba4a8c4, 0xbfc10ce59ba4a8c4,
+ 0x0, 0x8000000000000000, 0x3fe0000000000000, 0xbfe0000000000000, 0x3fc5555555555555, 0xbfc5555555555555, 0x3fc6d1cc1f0b5911, 0xbfc6d1cc1f0b5911,
+ 0x3fb49719f07a5377, 0xbfb49719f07a5377, 0x3fb5e875988eff8a, 0xbfb5e875988eff8a, 0x3fa46cdde5cc9718, 0xbfa46cdde5cc9718, 0x3fa5b58f3189ab72, 0xbfa5b58f3189ab72,
+ 0x0, 0x8000000000000000, 0x3fec34366179d426, 0xbfec34366179d426, 0x3fdecc2caec5160a, 0xbfdecc2caec5160a, 0x3fe0499e40c65ff4, 0xbfe0499e40c65ff4,
+ 0x3fcfacfb2399e637, 0xbfcfacfb2399e637, 0x3fd0ce61f7c6dc8c, 0xbfd0ce61f7c6dc8c, 0x3fbfead0b6996973, 0xbfbfead0b6996973, 0x3fc0f34e6531b848, 0xbfc0f34e6531b848,
+ 0x0, 0x8000000000000000, 0x3fe921fb54442d18, 0xbfe921fb54442d18, 0x3fddac670561bb4f, 0xbfddac670561bb4f, 0x3fdf40dd0b541417, 0xbfdf40dd0b541417,
+ 0x3fcf5b75f92c80dd, 0xbfcf5b75f92c80dd, 0x3fd09dc597d86362, 0xbfd09dc597d86362, 0x3fbfd5ba9aac2f6e, 0xbfbfd5ba9aac2f6e, 0x3fc0e6adccf40881, 0xbfc0e6adccf40881,
+ 0x0, 0x8000000000000000, 0x3fe921fb54442d18, 0xbfe921fb54442d18, 0x3fddac670561bb50, 0xbfddac670561bb50, 0x3fdf40dd0b541419, 0xbfdf40dd0b541419,
+ 0x3fcf5b75f92c80dd, 0xbfcf5b75f92c80dd, 0x3fd09dc597d86362, 0xbfd09dc597d86362, 0x3fbfd5ba9aac2f6d, 0xbfbfd5ba9aac2f6d, 0x3fc0e6adccf40881, 0xbfc0e6adccf40881,
+ 0x0, 0x0, 0x7ff0000000000000, 0xfff0000000000000, 0x3fe193ea7aad030a, 0xbfe193ea7aad030a, 0x3fe2f0b8e9ca2470, 0xbfe2f0b8e9ca2470,
+ 0x3fd058aefa811451, 0xbfd058aefa811451, 0x3fd16aeb1724557b, 0xbfd16aeb1724557b, 0x3fc015891c9eaef7, 0xbfc015891c9eaef7, 0x3fc119dce19bdbaf, 0xbfc119dce19bdbaf,
+ 0x0, 0x8000000000000000, 0x3fd0000000000000, 0xbfd0000000000000, 0x3fc2e4051d9df309, 0xbfc2e4051d9df309, 0x3fc3e58203d3c35a, 0xbfc3e58203d3c35a,
+ 0x3fb3f670b6bdc73d, 0xbfb3f670b6bdc73d, 0x3fb528120e5938eb, 0xbfb528120e5938eb, 0x3fa4444750777669, 0xbfa4444750777669, 0x3fa584e61024002a, 0xbfa584e61024002a,
+ 0x0, 0x8000000000000000, 0x3fd0000000000000, 0xbfd0000000000000, 0x3fc2e4051d9df309, 0xbfc2e4051d9df309, 0x3fc3e58203d3c35a, 0xbfc3e58203d3c35a,
+ 0x3fb3f670b6bdc73d, 0xbfb3f670b6bdc73d, 0x3fb528120e5938eb, 0xbfb528120e5938eb, 0x3fa4444750777668, 0xbfa4444750777668, 0x3fa584e610240029, 0xbfa584e610240029,
+ 0x3ff0000000000000, 0x3ff0000000000000, 0x3fe14a280fb5068b, 0x3fe14a280fb5068b, 0x3fec1528065b7d4e, 0x3fec1528065b7d4e, 0x3feb96eeef58840d, 0x3feb96eeef58840d,
+ 0x3fef01549f7deea3, 0x3fef01549f7deea3, 0x3feee0b1fbc0f11d, 0x3feee0b1fbc0f11d, 0x3fefc015527d5bd2, 0x3fefc015527d5bd2, 0x3fefb7db2bfe0694, 0x3fefb7db2bfe0694,
+ 0x3ff0000000000000, 0x3ff0000000000000, 0x3ff8b07551d9f551, 0x3ff8b07551d9f551, 0x3ff20ac1862ae8d1, 0x3ff20ac1862ae8d1, 0x3ff24fb8f22373db, 0x3ff24fb8f22373db,
+ 0x3ff080ab05ca6146, 0x3ff080ab05ca6146, 0x3ff0915a03c95705, 0x3ff0915a03c95705, 0x3ff0200aac16db6f, 0x3ff0200aac16db6f, 0x3ff0242d9a16a685, 0x3ff0242d9a16a685,
+ 0x3ff0000000000000, 0x3ff0000000000000, 0xbff0000000000000, 0xbff0000000000000, 0x0, 0x0, 0xbfb917a6bc29b42d, 0xbfb917a6bc29b42d,
+ 0x3fe6a09e667f3bce, 0x3fe6a09e667f3bce, 0x3fe57d69348ceca0, 0x3fe57d69348ceca0, 0x3fed906bcf328d46, 0x3fed906bcf328d46, 0x3fed4134d14dc939, 0x3fed4134d14dc939,
+ 0x0, 0x8000000000000000, 0x3feaed548f090cee, 0xbfeaed548f090cee, 0x3fdeaee8744b05ef, 0xbfdeaee8744b05ef, 0x3fe0362939c69954, 0xbfe0362939c69954,
+ 0x3fcfaaeed4f31578, 0xbfcfaaeed4f31578, 0x3fd0cd00cef36436, 0xbfd0cd00cef36436, 0x3fbfeaaeee86ee36, 0xbfbfeaaeee86ee36, 0x3fc0f3378ddd71d2, 0xbfc0f3378ddd71d2,
+ 0x0, 0x8000000000000000, 0x3ff2cd9fc44eb982, 0xbff2cd9fc44eb982, 0x3fe0acd00fe63b97, 0xbfe0acd00fe63b97, 0x3fe1cf9dd4bb0b07, 0xbfe1cf9dd4bb0b07,
+ 0x3fd02accd9d08102, 0xbfd02accd9d08102, 0x3fd1335ba15f1d6c, 0xbfd1335ba15f1d6c, 0x3fc00aaccd00d2f1, 0xbfc00aaccd00d2f1, 0x3fc10cce392720b0, 0xbfc10cce392720b0,
+ 0x0, 0x8000000000000000, 0x0, 0x8000000000000000, 0x3fefffffffffffff, 0xbfefffffffffffff, 0x3fefd88da3d12525, 0xbfefd88da3d12525,
+ 0x3fe6a09e667f3bcd, 0xbfe6a09e667f3bcd, 0x3fe7b5df226aafb0, 0xbfe7b5df226aafb0, 0x3fd87de2a6aea963, 0xbfd87de2a6aea963, 0x3fd9ef7943a8ed8b, 0xbfd9ef7943a8ed8b,
+ 0x0, 0x8000000000000000, 0x3ff8eb245cbee3a5, 0xbff8eb245cbee3a5, 0x3fe17b4f5bf3474a, 0xbfe17b4f5bf3474a, 0x3fe2cd98fea0ab88, 0xbfe2cd98fea0ab88,
+ 0x3fd05785a43c4c56, 0xbfd05785a43c4c56, 0x3fd16953ea9fb257, 0xbfd16953ea9fb257, 0x3fc01577af1511a5, 0xbfc01577af1511a5, 0x3fc119c537e4dc4d, 0xbfc119c537e4dc4d,
+ 0x0, 0x8000000000000000, 0x3fe85efab514f394, 0xbfe85efab514f394, 0x3fdd9353d7568af3, 0xbfdd9353d7568af3, 0x3fdf20211a36475d, 0xbfdf20211a36475d,
+ 0x3fcf597ea69a1c84, 0xbfcf597ea69a1c84, 0x3fd09c744f539be4, 0xbfd09c744f539be4, 0x3fbfd5992bc4b835, 0xbfbfd5992bc4b835, 0x3fc0e6973998da48, 0xbfc0e6973998da48,
+ 0x0, 0x0, 0x0, 0x0, 0x7ff0000000000000, 0xfff0000000000000, 0xc0244e6c595afdcd, 0x40244e6c595afdcd,
+ 0x3ff0000000000000, 0xbff0000000000000, 0x3ff1a73d55278c4b, 0xbff1a73d55278c4b, 0x3fda827999fcef32, 0xbfda827999fcef32, 0x3fdc5e87185e67b6, 0xbfdc5e87185e67b6,
+};
+
+double *input_data = argX;
+const int count_input = 16;
+const int max_function = 20;
+
+static void builtin_double_trigonometric_function(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};
+
+ if(cl_check_double() == false) return;
+
+ OCL_CREATE_KERNEL("builtin_double_trigonometric_function");
+
+ 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(long) * 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]) < 7);
+ if(abs(gpu_data[k] - expectResult[k]) > 7)
+ {
+ 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_trigonometric_function)
+}
--
2.7.4
More information about the Beignet
mailing list