diff options
author | Zhigang Gong <zhigang.gong@intel.com> | 2014-06-06 10:37:19 +0800 |
---|---|---|
committer | Zhigang Gong <zhigang.gong@intel.com> | 2014-06-09 11:50:20 +0800 |
commit | ae0c8737f8d1c192ba8786606638194a86945178 (patch) | |
tree | ae4890c23a64a3d6dd265453429369bb5eae9607 | |
parent | 71492287f76673fafce73a8aa6524db7c4a39342 (diff) |
utests: add a double precision check test case.
v2:
fix some bugs in test case.
Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
-rw-r--r-- | kernels/double_precision_check.cl | 11 | ||||
-rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
-rw-r--r-- | utests/compiler_double_precision.cpp | 43 |
3 files changed, 55 insertions, 0 deletions
diff --git a/kernels/double_precision_check.cl b/kernels/double_precision_check.cl new file mode 100644 index 0000000..e55cafa --- /dev/null +++ b/kernels/double_precision_check.cl @@ -0,0 +1,11 @@ +#pragma OPENCL EXTENSION cl_khr_fp64 : enable + +__kernel void +double_precision_check(__global float* src, __global float* dst) +{ + int id = (int)get_global_id(0); + double d0 = 0.12345678912345678 + src[1]; + double d1 = 0.12355678922345678 + src[0]; + float rem = d1 - d0; + dst[id] = rem; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 9ad08c9..1c523cb 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -167,6 +167,7 @@ set (utests_sources compiler_getelementptr_bitcast.cpp compiler_simd_any.cpp compiler_simd_all.cpp + compiler_double_precision.cpp load_program_from_bin.cpp profiling_exec.cpp enqueue_copy_buf.cpp diff --git a/utests/compiler_double_precision.cpp b/utests/compiler_double_precision.cpp new file mode 100644 index 0000000..217fd18 --- /dev/null +++ b/utests/compiler_double_precision.cpp @@ -0,0 +1,43 @@ +#include "utest_helper.hpp" +#include <math.h> + +static void double_precision_check(void) +{ + const size_t n = 16; //8192 * 4; + + double d0 = 0.12345678912345678; + double d1 = 0.12355678922345678; + float cpu_result = d1 - d0; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("double_precision_check"); + //OCL_CREATE_KERNEL("compiler_array"); + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); + for (uint32_t i = 0; i < n; ++i) ((float*)buf_data[0])[i] = 0; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + bool precisionOK = true; + for (uint32_t i = 0; i < n; ++i) { + float error = ((float*)buf_data[1])[i] - cpu_result; + if (error != 0) + precisionOK = false; + OCL_ASSERT((fabs(error) < 1e-4)); + } + if (!precisionOK) + printf("\n - WARN: GPU doesn't have correct double precision. Got %.7G, expected %.7G\n", ((float*)buf_data[1])[0], cpu_result); +} + +MAKE_UTEST_FROM_FUNCTION(double_precision_check); |