[Beignet] [PATCH] improve built-in function "sinpi"
Zhigang Gong
zhigang.gong at linux.intel.com
Fri Aug 30 00:09:34 PDT 2013
I tried this patchset and got the following compilation warnings.
Could you check and fix it in next version? Thanks.
Building CXX object utests/CMakeFiles/utests.dir/builtin_lgamma_r.cpp.o
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp: In function ‘float sinpi(float)’:
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp:19:20: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp:36:27: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp: In function ‘void builtin_sinpi()’:
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp:42:3: warning: ‘n’ may be used uninitialized in this function [-Wuninitialized]
/home/gongzg/git/fdo/beignet/utests/builtin_sinpi.cpp:16:7: note: ‘n’ was declared here
On Fri, Aug 23, 2013 at 02:35:34PM +0800, Homer Hsing wrote:
> "sinpi" was calculated as "sin(pi * x)".
> But that was not a quite-good way.
> This patch improved the function, also included a test case.
>
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
> backend/src/ocl_stdlib.tmpl.h | 56 ++++++++++++++++++++++++-
> kernels/builtin_sinpi.cl | 4 ++
> utests/CMakeLists.txt | 1 +
> utests/builtin_sinpi.cpp | 98 +++++++++++++++++++++++++++++++++++++++++++
> 4 files changed, 158 insertions(+), 1 deletion(-)
> create mode 100644 kernels/builtin_sinpi.cl
> create mode 100644 utests/builtin_sinpi.cpp
>
> diff --git a/backend/src/ocl_stdlib.tmpl.h b/backend/src/ocl_stdlib.tmpl.h
> index c428fac..a256a8d 100644
> --- a/backend/src/ocl_stdlib.tmpl.h
> +++ b/backend/src/ocl_stdlib.tmpl.h
> @@ -579,7 +579,61 @@ INLINE_OVERLOADABLE float __gen_ocl_internal_cospi(float x) {
> }
> INLINE_OVERLOADABLE float native_sin(float x) { return __gen_ocl_sin(x); }
> INLINE_OVERLOADABLE float __gen_ocl_internal_sinpi(float x) {
> - return __gen_ocl_sin(x * M_PI_F);
> +/*
> + * ====================================================
> + * Copyright (C) 1993 by Sun Microsystems, Inc. All rights reserved.
> + *
> + * Developed at SunPro, a Sun Microsystems, Inc. business.
> + * Permission to use, copy, modify, and distribute this
> + * software is freely granted, provided that this notice
> + * is preserved.
> + * ====================================================
> + */
> + float y, z;
> + int n, ix;
> + ix = *(int *) (&x) & 0x7fffffff;
> + if (ix < 0x3e800000)
> + return __gen_ocl_sin(M_PI_F * x);
> + y = -x;
> + z = __gen_ocl_rndd(y);
> + if (z != y) {
> + y *= 0.5f;
> + y = 2.f * (y - __gen_ocl_rndd(y));
> + n = y * 4.f;
> + } else {
> + if (ix >= 0x4b800000) {
> + y = 0;
> + n = 0;
> + } else {
> + if (ix < 0x4b000000)
> + z = y + 8.3886080000e+06f;
> + int n = *(int *) (&z);
> + n &= 1;
> + y = n;
> + n <<= 2;
> + }
> + }
> + switch (n) {
> + case 0:
> + y = __gen_ocl_sin(M_PI_F * y);
> + break;
> + case 1:
> + case 2:
> + y = __gen_ocl_cos(M_PI_F * (0.5f - y));
> + break;
> + case 3:
> + case 4:
> + y = __gen_ocl_sin(M_PI_F * (1.f - y));
> + break;
> + case 5:
> + case 6:
> + y = -__gen_ocl_cos(M_PI_F * (y - 1.5f));
> + break;
> + default:
> + y = __gen_ocl_sin(M_PI_F * (y - 2.f));
> + break;
> + }
> + return -y;
> }
> INLINE_OVERLOADABLE float native_sqrt(float x) { return __gen_ocl_sqrt(x); }
> INLINE_OVERLOADABLE float native_rsqrt(float x) { return __gen_ocl_rsqrt(x); }
> diff --git a/kernels/builtin_sinpi.cl b/kernels/builtin_sinpi.cl
> new file mode 100644
> index 0000000..134152d
> --- /dev/null
> +++ b/kernels/builtin_sinpi.cl
> @@ -0,0 +1,4 @@
> +kernel void builtin_sinpi(global float *src, global float *dst) {
> + int i = get_global_id(0);
> + dst[i] = sinpi(src[i]);
> +};
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 1aade66..31f9fbc 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -113,6 +113,7 @@ set (utests_sources
> builtin_shuffle.cpp
> builtin_shuffle2.cpp
> builtin_sign.cpp
> + builtin_sinpi.cpp
> buildin_work_dim.cpp
> builtin_global_size.cpp
> builtin_local_size.cpp
> diff --git a/utests/builtin_sinpi.cpp b/utests/builtin_sinpi.cpp
> new file mode 100644
> index 0000000..3a9bc50
> --- /dev/null
> +++ b/utests/builtin_sinpi.cpp
> @@ -0,0 +1,98 @@
> +#include <cmath>
> +#include "utest_helper.hpp"
> +
> +static float sinpi(float x) {
> +/*
> + * ====================================================
> + * Copyright (C) 1993 by Sun Microsystems, Inc. All rights reserved.
> + *
> + * Developed at SunPro, a Sun Microsystems, Inc. business.
> + * Permission to use, copy, modify, and distribute this
> + * software is freely granted, provided that this notice
> + * is preserved.
> + * ====================================================
> + */
> + float y, z;
> + int n, ix;
> + const float pi = 3.1415927410e+00f;
> +
> + ix = *(int *) (&x) & 0x7fffffff;
> +
> + if (ix < 0x3e800000)
> + return sinf(pi * x);
> + y = -x;
> + z = floorf(y);
> + if (z != y) {
> + y *= 0.5f;
> + y = 2.f * (y - floorf(y));
> + n = y * 4.f;
> + } else {
> + if (ix >= 0x4b800000) {
> + y = 0;
> + n = 0;
> + } else {
> + if (ix < 0x4b000000)
> + z = y + 8.3886080000e+06f;
> + int n = *(int *) (&z);
> + n &= 1;
> + y = n;
> + n <<= 2;
> + }
> + }
> + switch (n) {
> + case 0:
> + y = sinf(pi * y);
> + break;
> + case 1:
> + case 2:
> + y = cosf(pi * ((float) 0.5 - y));
> + break;
> + case 3:
> + case 4:
> + y = sinf(pi * (1.f - y));
> + break;
> + case 5:
> + case 6:
> + y = -cosf(pi * (y - (float) 1.5));
> + break;
> + default:
> + y = sinf(pi * (y - (float) 2.0));
> + break;
> + }
> + return -y;
> +}
> +
> +void builtin_sinpi(void)
> +{
> + const int n = 1024;
> + float src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("builtin_sinpi");
> + 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;
> +
> + for (int j = 0; j < 1000; j ++) {
> + OCL_MAP_BUFFER(0);
> + for (int i = 0; i < n; ++i) {
> + src[i] = ((float*)buf_data[0])[i] = (j*n + i) * 0.01f;
> + }
> + OCL_UNMAP_BUFFER(0);
> +
> + OCL_NDRANGE(1);
> +
> + OCL_MAP_BUFFER(1);
> + float *dst = (float*)buf_data[1];
> + for (int i = 0; i < n; ++i) {
> + float cpu = sinpi(src[i]);
> + OCL_ASSERT (fabsf(cpu - dst[i]) < 1e-4);
> + }
> + OCL_UNMAP_BUFFER(1);
> + }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(builtin_sinpi);
> --
> 1.8.1.2
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list