include fabs, floor, ceil, fract. Commit them
first because they are basic and related
Signed-off-by: rander <rander.w...@intel.com>
---
kernels/builtin_double_fabs_floor_ceil_fract.cl | 23 +++++++
utests/CMakeLists.txt | 3 +-
utests/builtin_double_fabs_floor_ceil_fract.cpp | 89 +++++++++++++++++++++++++
3 files changed, 114 insertions(+), 1 deletion(-)
create mode 100644 kernels/builtin_double_fabs_floor_ceil_fract.cl
create mode 100644 utests/builtin_double_fabs_floor_ceil_fract.cpp
diff --git a/kernels/builtin_double_fabs_floor_ceil_fract.cl
b/kernels/builtin_double_fabs_floor_ceil_fract.cl
new file mode 100644
index 0000000..a4b6ff1
--- /dev/null
+++ b/kernels/builtin_double_fabs_floor_ceil_fract.cl
@@ -0,0 +1,23 @@
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void builtin_double_fabs_floor_ceil_fract(__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++] = fabs(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = floor(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = ceil(X[j]);
+
+ for(j = 0; j < max_input; j++)
+ Z[i++] = fract(X[j], &dfloor);
+}
+
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 0f69a0e..10b0172 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -307,7 +307,8 @@ set (utests_sources
builtin_convert_int8toDouble.cpp
builtin_convert_int16toDouble.cpp
builtin_convert_double2float.cpp
- builtin_convert_ulong2double.cpp)
+ builtin_convert_ulong2double.cpp
+ builtin_double_fabs_floor_ceil_fract.cpp)
if (LLVM_VERSION_NODOT VERSION_GREATER 34)
SET(utests_sources
diff --git a/utests/builtin_double_fabs_floor_ceil_fract.cpp
b/utests/builtin_double_fabs_floor_ceil_fract.cpp
new file mode 100644
index 0000000..0aab6d2
--- /dev/null
+++ b/utests/builtin_double_fabs_floor_ceil_fract.cpp
@@ -0,0 +1,89 @@
+#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 fabs(double x)"
+ "OVERLOADABLE double floor(double x)"
+ "OVERLOADABLE double ceil(double x)"
+ "OVERLOADABLE double fract(double x, global double *p)"
+};
+
+unsigned long expectResult[] {
+ 0x0, 0x0, 0x2e30000000000000, 0x1000, 0x4000000001000000,
0x401ffffff0000000, 0x41f01fffffffff00, 0x42b0ffffffff0000,
+ 0x43f7ffffff000000, 0x45bc000000000000, 0x4ff000000000f000,
0x5ff0000000000000, 0x6ff0000000000000, 0x7fe0000000000001, 0x7fe0000000000000,
0x0,
+ 0x0, 0x0, 0x0, 0xbff0000000000000, 0x4000000000000000, 0xc020000000000000,
0x41f01ffffff00000, 0xc2b0ffffffff0000,
+ 0x43f7ffffff000000, 0xc5bc000000000000, 0xcff000000000f000,
0x5ff0000000000000, 0xeff0000000000000, 0x7fe0000000000001, 0xffe0000000000000,
0x0,
+ 0x0, 0x0, 0x3ff0000000000000, 0x0, 0x4008000000000000, 0xc01c000000000000,
0x41f0200000000000, 0xc2b0ffffffff0000,
+ 0x43f7ffffff000000, 0xc5bc000000000000, 0xcff000000000f000,
0x5ff0000000000000, 0xeff0000000000000, 0x7fe0000000000001, 0xffe0000000000000,
0x0,
+ 0x0, 0x8000000000000000, 0x2e30000000000000, 0x3fefffffffffffff,
0x3e40000000000000, 0x3e90000000000000, 0x3feffe0000000000, 0x0,
+ 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0
+};
+
+double *input_data = argX;
+const int count_input = 16;
+const int max_function = 4;
+
+static void builtin_double_fabs_floor_ceil_fract(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_fabs_floor_ceil_fract");
+
+ 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(gpu_data[k] == expectResult[k]);
+ if(gpu_data[k] != expectResult[k])
+ {
+ 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_fabs_floor_ceil_fract)
+}
--
2.7.4
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
