<div dir="ltr"> <div class="gmail_extra"> <div class="gmail_quote">On Wed, Jun 19, 2013 at 9:09 AM, Xing, Homer <<a href="mailto:homer.xing@intel.com" target="_blank">homer.xing@intel.com</a>> wrote: <blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Why do you "#define FLOAT double"? Why not directly use "double"? </blockquote><div style> I already give the reason in my cover letter. You can just change the macro to verify the same logic with float or double.</div> <div style> As it always fails with double, I need to verify the test case's correctness with float type. Any other question? </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> <div><div class="h5"> -----Original Message----- From: Zhigang Gong [mailto:<a href="mailto:zhigang.gong@linux.intel.com">zhigang.gong@linux.intel.com</a>] Sent: Tuesday, June 18, 2013 5:29 PM To: Xing, Homer Cc: <a href="mailto:beignet@lists.freedesktop.org">beignet@lists.freedesktop.org</a> 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 <<a href="mailto:homer.xing@intel.com">homer.xing@intel.com</a>> > > Signed-off-by: Homer Hsing <<a href="mailto:homer.xing@intel.com">homer.xing@intel.com</a>> > Signed-off-by: Zhigang Gong <<a href="mailto:zhigang.gong@linux.intel.com">zhigang.gong@linux.intel.com</a>> > --- > kernels/<a href="http://compiler_double.cl" target="_blank">compiler_double.cl</a> | 10 +++++++++ > kernels/<a href="http://compiler_double_2.cl" target="_blank">compiler_double_2.cl</a> | 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/<a href="http://compiler_double.cl" target="_blank">compiler_double.cl</a> create mode 100644 > kernels/<a href="http://compiler_double_2.cl" target="_blank">compiler_double_2.cl</a> create mode 100644 > utests/compiler_double.cpp create mode 100644 > utests/compiler_double_2.cpp > > diff --git a/kernels/<a href="http://compiler_double.cl" target="_blank">compiler_double.cl</a> b/kernels/<a href="http://compiler_double.cl" target="_blank">compiler_double.cl</a> > new file mode 100644 index 0000000..6d00d8a > --- /dev/null > +++ b/kernels/<a href="http://compiler_double.cl" target="_blank">compiler_double.cl</a> > @@ -0,0 +1,10 @@ </div></div>> +#pragma OPENCL EXTENSION cl_khr_fp64 : enable #define FLOAT double <div class="im">> +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/<a href="http://compiler_double_2.cl" target="_blank">compiler_double_2.cl</a> > b/kernels/<a href="http://compiler_double_2.cl" target="_blank">compiler_double_2.cl</a> new file mode 100644 index > 0000000..9e5c5ec > --- /dev/null > +++ b/kernels/<a href="http://compiler_double_2.cl" target="_blank">compiler_double_2.cl</a> > @@ -0,0 +1,6 @@ </div>> +#pragma OPENCL EXTENSION cl_khr_fp64 : enable kernel void > +compiler_double_2(global float *src, global double *dst) { <div><div class="h5">> + 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); </div></div>> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); OCL_SET_ARG(1, > + sizeof(cl_mem), &buf[1]); globals[0] = n; locals[0] = 16; <div><div class="h5">> + > + // 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); </div></div>> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); OCL_SET_ARG(1, > + sizeof(cl_mem), &buf[1]); globals[0] = n; locals[0] = 16; <div class="HOEnZb"><div class="h5">> + > + // 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 > <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a> > <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a> _______________________________________________ Beignet mailing list <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a> <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a> </div></div></blockquote></div> </div></div>