[Beignet] [PATCH 2/2] utests: add a double precision check test case.

[Zhigang Gong]

On Fri, Jun 06, 2014 at 10:53:35AM +0800, Zhigang Gong wrote:
> Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
> ---
>  kernels/double_precision_check.cl    | 11 +++++++++
>  utests/CMakeLists.txt                |  1 +
>  utests/compiler_double_precision.cpp | 43 ++++++++++++++++++++++++++++++++++++
>  3 files changed, 55 insertions(+)
>  create mode 100644 kernels/double_precision_check.cl
>  create mode 100644 utests/compiler_double_precision.cpp
> 
> diff --git a/kernels/double_precision_check.cl b/kernels/double_precision_check.cl
> new file mode 100644
> index 0000000..6dd1123
> --- /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.12355678912345688 + src[0];
d1 should be the same as on cpu side which is 0.12355678922345678
> +  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..d422132
> --- /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 - cpu_result != 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);
> -- 
> 1.8.3.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet