[Beignet] [PATCH 2/2] test case for 64-bit float
zhigang gong
zhigang.gong at gmail.com
Tue Jun 18 18:36:28 PDT 2013
On Wed, Jun 19, 2013 at 9:09 AM, Xing, Homer <homer.xing at intel.com> wrote:
> Why do you "#define FLOAT double"? Why not directly use "double"?
>
I already give the reason in my cover letter. You can just change the
macro to verify the same logic with float or double.
As it always fails with double, I need to verify the test case's
correctness with float type. Any other question?
>
> -----Original Message-----
> From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
> Sent: Tuesday, June 18, 2013 5:29 PM
> To: Xing, Homer
> Cc: beignet at lists.freedesktop.org
> Subject: Re: [Beignet] [PATCH 2/2] test case for 64-bit float
>
> 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
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130619/8edca9ff/attachment.html>
More information about the Beignet
mailing list