[Beignet] [PATCH 2/2] test case for 64-bit float

[Zhigang Gong]

On Tue, Jun 18, 2013 at 05:10:46PM +0800, Zhigang Gong wrote:
> From: Homer Hsing <homer.xing at intel.com>
> 
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
> ---
>  kernels/compiler_double.cl   | 10 +++++++++
>  kernels/compiler_double_2.cl |  6 ++++++
>  utests/CMakeLists.txt        |  2 ++
>  utests/compiler_double.cpp   | 51 ++++++++++++++++++++++++++++++++++++++++++++
>  utests/compiler_double_2.cpp | 47 ++++++++++++++++++++++++++++++++++++++++
>  5 files changed, 116 insertions(+)
>  create mode 100644 kernels/compiler_double.cl
>  create mode 100644 kernels/compiler_double_2.cl
>  create mode 100644 utests/compiler_double.cpp
>  create mode 100644 utests/compiler_double_2.cpp
> 
> diff --git a/kernels/compiler_double.cl b/kernels/compiler_double.cl
> new file mode 100644
> index 0000000..6d00d8a
> --- /dev/null
> +++ b/kernels/compiler_double.cl
> @@ -0,0 +1,10 @@
> +#pragma OPENCL EXTENSION cl_khr_fp64 : enable
> +#define FLOAT double
> +kernel void compiler_double(global FLOAT *src, global FLOAT *dst) {
> +  int i = get_global_id(0);
> +  FLOAT d = 1.234567890123456789;
> +  if (i % 2 == 0)
> +    dst[i] = d * (src[i] + d);
> +  else
> +    dst[i] = d * src[i];
> +}
> diff --git a/kernels/compiler_double_2.cl b/kernels/compiler_double_2.cl
> new file mode 100644
> index 0000000..9e5c5ec
> --- /dev/null
> +++ b/kernels/compiler_double_2.cl
> @@ -0,0 +1,6 @@
> +#pragma OPENCL EXTENSION cl_khr_fp64 : enable
> +kernel void compiler_double_2(global float *src, global double *dst) {
> +  int i = get_global_id(0);
> +  float d = 1.234567890123456789f;
> +  dst[i] = d * (d + src[i]);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 93778ed..0f740d2 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -27,6 +27,8 @@ set (utests_sources
>    compiler_copy_image.cpp
>    compiler_copy_image_3d.cpp
>    compiler_copy_buffer_row.cpp
> +  compiler_double.cpp
> +  compiler_double_2.cpp
>    compiler_fabs.cpp
>    compiler_fill_image.cpp
>    compiler_fill_image0.cpp
> diff --git a/utests/compiler_double.cpp b/utests/compiler_double.cpp
> new file mode 100644
> index 0000000..d92f264
> --- /dev/null
> +++ b/utests/compiler_double.cpp
> @@ -0,0 +1,51 @@
> +#include <cmath>
> +#include "utest_helper.hpp"
> +
> +#define FLOAT double
> +
> +static void cpu(int global_id, FLOAT *src, FLOAT *dst) {
> +  FLOAT f = src[global_id];
> +  FLOAT d = 1.234567890123456789;
> +  if (global_id % 2 == 0)
> +    dst[global_id] = d * (f + d);
> +  else
> +    dst[global_id] = d * f;
> +}
> +
> +void compiler_FLOAT(void)
> +{
> +  const size_t n = 16;
> +  FLOAT cpu_dst[n], cpu_src[n];
> +
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL("compiler_double");
> +  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(FLOAT), NULL);
> +  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(FLOAT), NULL);
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> +  globals[0] = n;
> +  locals[0] = 16;
> +
> +  // Run random tests
> +  for (uint32_t pass = 0; pass < 1; ++pass) {
> +    OCL_MAP_BUFFER(0);
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      cpu_src[i] = ((FLOAT*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
> +    OCL_UNMAP_BUFFER(0);
> +
> +    // Run the kernel on GPU
> +    OCL_NDRANGE(1);
> +
> +    // Run on CPU
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      cpu(i, cpu_src, cpu_dst);
> +
> +    // Compare
> +    OCL_MAP_BUFFER(1);
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      OCL_ASSERT(fabs(((FLOAT*)buf_data[1])[i] - cpu_dst[i]) < 1e-4);
> +    OCL_UNMAP_BUFFER(1);
> +  }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_FLOAT);
                 sorry, one typo here, should be compiler_double rather than compiler_FLOAT.
> diff --git a/utests/compiler_double_2.cpp b/utests/compiler_double_2.cpp
> new file mode 100644
> index 0000000..8c30443
> --- /dev/null
> +++ b/utests/compiler_double_2.cpp
> @@ -0,0 +1,47 @@
> +#include <cmath>
> +#include "utest_helper.hpp"
> +
> +static void cpu(int global_id, float *src, double *dst) {
> +  float f = src[global_id];
> +  float d = 1.234567890123456789;
> +  dst[global_id] = d * (d + f);
> +}
> +
> +void compiler_double_2(void)
> +{
> +  const size_t n = 16;
> +  float cpu_src[n];
> +  double cpu_dst[n];
> +
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL("compiler_double_2");
> +  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
> +  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(double), NULL);
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> +  globals[0] = n;
> +  locals[0] = 16;
> +
> +  // Run random tests
> +  for (uint32_t pass = 0; pass < 1; ++pass) {
> +    OCL_MAP_BUFFER(0);
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      cpu_src[i] = ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
> +    OCL_UNMAP_BUFFER(0);
> +
> +    // Run the kernel on GPU
> +    OCL_NDRANGE(1);
> +
> +    // Run on CPU
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      cpu(i, cpu_src, cpu_dst);
> +
> +    // Compare
> +    OCL_MAP_BUFFER(1);
> +    for (int32_t i = 0; i < (int32_t) n; ++i)
> +      OCL_ASSERT(fabs(((double*)buf_data[1])[i] - cpu_dst[i]) < 1e-4);
> +    OCL_UNMAP_BUFFER(1);
> +  }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_double_2);
> -- 
> 1.7.11.7
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet