[Beignet] [PATCH] add built-in function "lgamma", "lgamma_r"
Zhigang Gong
zhigang.gong at linux.intel.com
Fri Aug 30 00:14:20 PDT 2013
LGTM, pushed, thanks.
On Mon, Aug 26, 2013 at 12:51:53PM +0800, Homer Hsing wrote:
> also include test cases
>
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
> backend/src/builtin_vector_proto.def | 14 +-
> backend/src/ocl_stdlib.tmpl.h | 384 +++++++++++++++++++++++++++++++++++
> kernels/builtin_lgamma.cl | 4 +
> kernels/builtin_lgamma_r.cl | 4 +
> utests/CMakeLists.txt | 2 +
> utests/builtin_lgamma.cpp | 40 ++++
> utests/builtin_lgamma_r.cpp | 46 +++++
> 7 files changed, 487 insertions(+), 7 deletions(-)
> create mode 100644 kernels/builtin_lgamma.cl
> create mode 100644 kernels/builtin_lgamma_r.cl
> create mode 100644 utests/builtin_lgamma.cpp
> create mode 100644 utests/builtin_lgamma_r.cpp
>
> diff --git a/backend/src/builtin_vector_proto.def b/backend/src/builtin_vector_proto.def
> index 2a057bb..2a3daf2 100644
> --- a/backend/src/builtin_vector_proto.def
> +++ b/backend/src/builtin_vector_proto.def
> @@ -61,13 +61,13 @@ float ldexp (float x, int k)
> doublen ldexp (doublen x, intn k)
> doublen ldexp (doublen x, int k)
> double ldexp (double x, int k)
> -#gentype lgamma (gentype x)
> -#floatn lgamma_r (floatn x, __global intn *signp)
> -#floatn lgamma_r (floatn x, __local intn *signp)
> -#floatn lgamma_r (floatn x, __private intn *signp)
> -#float lgamma_r (float x, __global int *signp)
> -#float lgamma_r (float x, __local int *signp)
> -#float lgamma_r (float x, __private int *signp)
> +gentype lgamma (gentype x)
> +floatn lgamma_r (floatn x, __global intn *signp)
> +floatn lgamma_r (floatn x, __local intn *signp)
> +floatn lgamma_r (floatn x, __private intn *signp)
> +float lgamma_r (float x, __global int *signp)
> +float lgamma_r (float x, __local int *signp)
> +float lgamma_r (float x, __private int *signp)
> #doublen lgamma_r (doublen x, __global intn *signp)
> #doublen lgamma_r (doublen x, __local intn *signp)
> #doublen lgamma_r (doublen x, __private intn *signp)
> diff --git a/backend/src/ocl_stdlib.tmpl.h b/backend/src/ocl_stdlib.tmpl.h
> index c8d20b6..ac1999d 100644
> --- a/backend/src/ocl_stdlib.tmpl.h
> +++ b/backend/src/ocl_stdlib.tmpl.h
> @@ -823,6 +823,390 @@ INLINE_OVERLOADABLE float tgamma(float x) {
> r = nadj - r;
> return r;
> }
> +
> +INLINE_OVERLOADABLE float lgamma(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.
> + * ====================================================
> + */
> + const float
> + zero= 0.,
> + one = 1.0000000000e+00,
> + pi = 3.1415927410e+00,
> + a0 = 7.7215664089e-02,
> + a1 = 3.2246702909e-01,
> + a2 = 6.7352302372e-02,
> + a3 = 2.0580807701e-02,
> + a4 = 7.3855509982e-03,
> + a5 = 2.8905137442e-03,
> + a6 = 1.1927076848e-03,
> + a7 = 5.1006977446e-04,
> + a8 = 2.2086278477e-04,
> + a9 = 1.0801156895e-04,
> + a10 = 2.5214456400e-05,
> + a11 = 4.4864096708e-05,
> + tc = 1.4616321325e+00,
> + tf = -1.2148628384e-01,
> + tt = 6.6971006518e-09,
> + t0 = 4.8383611441e-01,
> + t1 = -1.4758771658e-01,
> + t2 = 6.4624942839e-02,
> + t3 = -3.2788541168e-02,
> + t4 = 1.7970675603e-02,
> + t5 = -1.0314224288e-02,
> + t6 = 6.1005386524e-03,
> + t7 = -3.6845202558e-03,
> + t8 = 2.2596477065e-03,
> + t9 = -1.4034647029e-03,
> + t10 = 8.8108185446e-04,
> + t11 = -5.3859531181e-04,
> + t12 = 3.1563205994e-04,
> + t13 = -3.1275415677e-04,
> + t14 = 3.3552918467e-04,
> + u0 = -7.7215664089e-02,
> + u1 = 6.3282704353e-01,
> + u2 = 1.4549225569e+00,
> + u3 = 9.7771751881e-01,
> + u4 = 2.2896373272e-01,
> + u5 = 1.3381091878e-02,
> + v1 = 2.4559779167e+00,
> + v2 = 2.1284897327e+00,
> + v3 = 7.6928514242e-01,
> + v4 = 1.0422264785e-01,
> + v5 = 3.2170924824e-03,
> + s0 = -7.7215664089e-02,
> + s1 = 2.1498242021e-01,
> + s2 = 3.2577878237e-01,
> + s3 = 1.4635047317e-01,
> + s4 = 2.6642270386e-02,
> + s5 = 1.8402845599e-03,
> + s6 = 3.1947532989e-05,
> + r1 = 1.3920053244e+00,
> + r2 = 7.2193557024e-01,
> + r3 = 1.7193385959e-01,
> + r4 = 1.8645919859e-02,
> + r5 = 7.7794247773e-04,
> + r6 = 7.3266842264e-06,
> + w0 = 4.1893854737e-01,
> + w1 = 8.3333335817e-02,
> + w2 = -2.7777778450e-03,
> + w3 = 7.9365057172e-04,
> + w4 = -5.9518753551e-04,
> + w5 = 8.3633989561e-04,
> + w6 = -1.6309292987e-03;
> + float t, y, z, nadj, p, p1, p2, p3, q, r, w;
> + int i, hx, ix;
> + nadj = 0;
> + hx = *(int *)&x;
> + ix = hx & 0x7fffffff;
> + if (ix >= 0x7f800000)
> + return x * x;
> + if (ix == 0)
> + return ((x + one) / zero);
> + if (ix < 0x1c800000) {
> + if (hx < 0) {
> + return -native_log(-x);
> + } else
> + return -native_log(x);
> + }
> + if (hx < 0) {
> + if (ix >= 0x4b000000)
> + return ((-x) / zero);
> + t = __gen_ocl_internal_sinpi(x);
> + if (t == zero)
> + return ((-x) / zero);
> + nadj = native_log(pi / __gen_ocl_fabs(t * x));
> + x = -x;
> + }
> + if (ix == 0x3f800000 || ix == 0x40000000)
> + r = 0;
> + else if (ix < 0x40000000) {
> + if (ix <= 0x3f666666) {
> + r = -native_log(x);
> + if (ix >= 0x3f3b4a20) {
> + y = one - x;
> + i = 0;
> + } else if (ix >= 0x3e6d3308) {
> + y = x - (tc - one);
> + i = 1;
> + } else {
> + y = x;
> + i = 2;
> + }
> + } else {
> + r = zero;
> + if (ix >= 0x3fdda618) {
> + y = (float) 2.0 - x;
> + i = 0;
> + }
> + else if (ix >= 0x3F9da620) {
> + y = x - tc;
> + i = 1;
> + }
> + else {
> + y = x - one;
> + i = 2;
> + }
> + }
> + switch (i) {
> + case 0:
> + z = y * y;
> + p1 = a0 + z * (a2 + z * (a4 + z * (a6 + z * (a8 + z * a10))));
> + p2 = z * (a1 + z * (a3 + z * (a5 + z * (a7 + z * (a9 + z * a11)))));
> + p = y * p1 + p2;
> + r += (p - (float) 0.5 * y);
> + break;
> + case 1:
> + z = y * y;
> + w = z * y;
> + p1 = t0 + w * (t3 + w * (t6 + w * (t9 + w * t12)));
> + p2 = t1 + w * (t4 + w * (t7 + w * (t10 + w * t13)));
> + p3 = t2 + w * (t5 + w * (t8 + w * (t11 + w * t14)));
> + p = z * p1 - (tt - w * (p2 + y * p3));
> + r += (tf + p);
> + break;
> + case 2:
> + p1 = y * (u0 + y * (u1 + y * (u2 + y * (u3 + y * (u4 + y * u5)))));
> + p2 = one + y * (v1 + y * (v2 + y * (v3 + y * (v4 + y * v5))));
> + r += (-(float) 0.5 * y + p1 / p2);
> + }
> + } else if (ix < 0x41000000) {
> + i = (int) x;
> + t = zero;
> + y = x - (float) i;
> + p = y * (s0 + y * (s1 + y * (s2 + y * (s3 + y * (s4 + y * (s5 + y * s6))))));
> + q = one + y * (r1 + y * (r2 + y * (r3 + y * (r4 + y * (r5 + y * r6)))));
> + r = .5f * y + p / q;
> + z = one;
> + switch (i) {
> + case 7:
> + z *= (y + (float) 6.0);
> + case 6:
> + z *= (y + (float) 5.0);
> + case 5:
> + z *= (y + (float) 4.0);
> + case 4:
> + z *= (y + (float) 3.0);
> + case 3:
> + z *= (y + (float) 2.0);
> + r += native_log(z);
> + break;
> + }
> +
> + } else if (ix < 0x5c800000) {
> + t = native_log(x);
> + z = one / x;
> + y = z * z;
> + w = w0 + z * (w1 + y * (w2 + y * (w3 + y * (w4 + y * (w5 + y * w6)))));
> + r = (x - .5f) * (t - one) + w;
> + } else
> + r = x * (native_log(x) - one);
> + if (hx < 0)
> + r = nadj - r;
> + return r;
> +}
> +
> +/*
> + * ====================================================
> + * 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.
> + * ====================================================
> + */
> +#define BODY \
> + const float \
> + zero= 0., \
> + one = 1.0000000000e+00, \
> + pi = 3.1415927410e+00, \
> + a0 = 7.7215664089e-02, \
> + a1 = 3.2246702909e-01, \
> + a2 = 6.7352302372e-02, \
> + a3 = 2.0580807701e-02, \
> + a4 = 7.3855509982e-03, \
> + a5 = 2.8905137442e-03, \
> + a6 = 1.1927076848e-03, \
> + a7 = 5.1006977446e-04, \
> + a8 = 2.2086278477e-04, \
> + a9 = 1.0801156895e-04, \
> + a10 = 2.5214456400e-05, \
> + a11 = 4.4864096708e-05, \
> + tc = 1.4616321325e+00, \
> + tf = -1.2148628384e-01, \
> + tt = 6.6971006518e-09, \
> + t0 = 4.8383611441e-01, \
> + t1 = -1.4758771658e-01, \
> + t2 = 6.4624942839e-02, \
> + t3 = -3.2788541168e-02, \
> + t4 = 1.7970675603e-02, \
> + t5 = -1.0314224288e-02, \
> + t6 = 6.1005386524e-03, \
> + t7 = -3.6845202558e-03, \
> + t8 = 2.2596477065e-03, \
> + t9 = -1.4034647029e-03, \
> + t10 = 8.8108185446e-04, \
> + t11 = -5.3859531181e-04, \
> + t12 = 3.1563205994e-04, \
> + t13 = -3.1275415677e-04, \
> + t14 = 3.3552918467e-04, \
> + u0 = -7.7215664089e-02, \
> + u1 = 6.3282704353e-01, \
> + u2 = 1.4549225569e+00, \
> + u3 = 9.7771751881e-01, \
> + u4 = 2.2896373272e-01, \
> + u5 = 1.3381091878e-02, \
> + v1 = 2.4559779167e+00, \
> + v2 = 2.1284897327e+00, \
> + v3 = 7.6928514242e-01, \
> + v4 = 1.0422264785e-01, \
> + v5 = 3.2170924824e-03, \
> + s0 = -7.7215664089e-02, \
> + s1 = 2.1498242021e-01, \
> + s2 = 3.2577878237e-01, \
> + s3 = 1.4635047317e-01, \
> + s4 = 2.6642270386e-02, \
> + s5 = 1.8402845599e-03, \
> + s6 = 3.1947532989e-05, \
> + r1 = 1.3920053244e+00, \
> + r2 = 7.2193557024e-01, \
> + r3 = 1.7193385959e-01, \
> + r4 = 1.8645919859e-02, \
> + r5 = 7.7794247773e-04, \
> + r6 = 7.3266842264e-06, \
> + w0 = 4.1893854737e-01, \
> + w1 = 8.3333335817e-02, \
> + w2 = -2.7777778450e-03, \
> + w3 = 7.9365057172e-04, \
> + w4 = -5.9518753551e-04, \
> + w5 = 8.3633989561e-04, \
> + w6 = -1.6309292987e-03; \
> + float t, y, z, nadj, p, p1, p2, p3, q, r, w; \
> + int i, hx, ix; \
> + nadj = 0; \
> + hx = *(int *)&x; \
> + *signgamp = 1; \
> + ix = hx & 0x7fffffff; \
> + if (ix >= 0x7f800000) \
> + return x * x; \
> + if (ix == 0) \
> + return ((x + one) / zero); \
> + if (ix < 0x1c800000) { \
> + if (hx < 0) { \
> + *signgamp = -1; \
> + return -native_log(-x); \
> + } else \
> + return -native_log(x); \
> + } \
> + if (hx < 0) { \
> + if (ix >= 0x4b000000) \
> + return ((-x) / zero); \
> + t = __gen_ocl_internal_sinpi(x); \
> + if (t == zero) \
> + return ((-x) / zero); \
> + nadj = native_log(pi / __gen_ocl_fabs(t * x)); \
> + if (t < zero) \
> + *signgamp = -1; \
> + x = -x; \
> + } \
> + if (ix == 0x3f800000 || ix == 0x40000000) \
> + r = 0; \
> + else if (ix < 0x40000000) { \
> + if (ix <= 0x3f666666) { \
> + r = -native_log(x); \
> + if (ix >= 0x3f3b4a20) { \
> + y = one - x; \
> + i = 0; \
> + } else if (ix >= 0x3e6d3308) { \
> + y = x - (tc - one); \
> + i = 1; \
> + } else { \
> + y = x; \
> + i = 2; \
> + } \
> + } else { \
> + r = zero; \
> + if (ix >= 0x3fdda618) { \
> + y = (float) 2.0 - x; \
> + i = 0; \
> + } \
> + else if (ix >= 0x3F9da620) { \
> + y = x - tc; \
> + i = 1; \
> + } \
> + else { \
> + y = x - one; \
> + i = 2; \
> + } \
> + } \
> + switch (i) { \
> + case 0: \
> + z = y * y; \
> + p1 = a0 + z * (a2 + z * (a4 + z * (a6 + z * (a8 + z * a10)))); \
> + p2 = z * (a1 + z * (a3 + z * (a5 + z * (a7 + z * (a9 + z * a11))))); \
> + p = y * p1 + p2; \
> + r += (p - (float) 0.5 * y); \
> + break; \
> + case 1: \
> + z = y * y; \
> + w = z * y; \
> + p1 = t0 + w * (t3 + w * (t6 + w * (t9 + w * t12))); \
> + p2 = t1 + w * (t4 + w * (t7 + w * (t10 + w * t13))); \
> + p3 = t2 + w * (t5 + w * (t8 + w * (t11 + w * t14))); \
> + p = z * p1 - (tt - w * (p2 + y * p3)); \
> + r += (tf + p); \
> + break; \
> + case 2: \
> + p1 = y * (u0 + y * (u1 + y * (u2 + y * (u3 + y * (u4 + y * u5))))); \
> + p2 = one + y * (v1 + y * (v2 + y * (v3 + y * (v4 + y * v5)))); \
> + r += (-(float) 0.5 * y + p1 / p2); \
> + } \
> + } else if (ix < 0x41000000) { \
> + i = (int) x; \
> + t = zero; \
> + y = x - (float) i; \
> + p = y * (s0 + y * (s1 + y * (s2 + y * (s3 + y * (s4 + y * (s5 + y * s6)))))); \
> + q = one + y * (r1 + y * (r2 + y * (r3 + y * (r4 + y * (r5 + y * r6))))); \
> + r = .5f * y + p / q; \
> + z = one; \
> + switch (i) { \
> + case 7: \
> + z *= (y + (float) 6.0); \
> + case 6: \
> + z *= (y + (float) 5.0); \
> + case 5: \
> + z *= (y + (float) 4.0); \
> + case 4: \
> + z *= (y + (float) 3.0); \
> + case 3: \
> + z *= (y + (float) 2.0); \
> + r += native_log(z); \
> + break; \
> + } \
> + \
> + } else if (ix < 0x5c800000) { \
> + t = native_log(x); \
> + z = one / x; \
> + y = z * z; \
> + w = w0 + z * (w1 + y * (w2 + y * (w3 + y * (w4 + y * (w5 + y * w6))))); \
> + r = (x - .5f) * (t - one) + w; \
> + } else \
> + r = x * (native_log(x) - one); \
> + if (hx < 0) \
> + r = nadj - r; \
> + return r;
> +INLINE_OVERLOADABLE float lgamma_r(float x, global int *signgamp) { BODY; }
> +INLINE_OVERLOADABLE float lgamma_r(float x, local int *signgamp) { BODY; }
> +INLINE_OVERLOADABLE float lgamma_r(float x, private int *signgamp) { BODY; }
> +#undef BODY
> +
> INLINE_OVERLOADABLE float native_log10(float x) {
> return native_log2(x) * 0.3010299956f;
> }
> diff --git a/kernels/builtin_lgamma.cl b/kernels/builtin_lgamma.cl
> new file mode 100644
> index 0000000..85bf859
> --- /dev/null
> +++ b/kernels/builtin_lgamma.cl
> @@ -0,0 +1,4 @@
> +kernel void builtin_lgamma(global float *src, global float *dst) {
> + int i = get_global_id(0);
> + dst[i] = lgamma(src[i]);
> +};
> diff --git a/kernels/builtin_lgamma_r.cl b/kernels/builtin_lgamma_r.cl
> new file mode 100644
> index 0000000..71fcc36
> --- /dev/null
> +++ b/kernels/builtin_lgamma_r.cl
> @@ -0,0 +1,4 @@
> +kernel void builtin_lgamma_r(global float *src, global float *dst, global int *signp) {
> + int i = get_global_id(0);
> + dst[i] = lgamma_r(src[i], signp+i);
> +};
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 949c6b7..08b4e32 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -114,6 +114,8 @@ set (utests_sources
> builtin_shuffle2.cpp
> builtin_sign.cpp
> builtin_sinpi.cpp
> + builtin_lgamma.cpp
> + builtin_lgamma_r.cpp
> builtin_tgamma.cpp
> buildin_work_dim.cpp
> builtin_global_size.cpp
> diff --git a/utests/builtin_lgamma.cpp b/utests/builtin_lgamma.cpp
> new file mode 100644
> index 0000000..876699a
> --- /dev/null
> +++ b/utests/builtin_lgamma.cpp
> @@ -0,0 +1,40 @@
> +#include <cmath>
> +#include "utest_helper.hpp"
> +
> +void builtin_lgamma(void) {
> + const int n = 1024;
> + float src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("builtin_lgamma");
> + 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 < 1024; j++) {
> + OCL_MAP_BUFFER(0);
> + for (int i = 0; i < n; ++i) {
> + src[i] = ((float*) buf_data[0])[i] = (j * n + i + 1) * 0.001f;
> + }
> + 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 = lgamma(src[i]);
> + float gpu = dst[i];
> + if (fabsf(cpu - gpu) >= 1e-3) {
> + printf("%f %f %f\n", src[i], cpu, gpu);
> + OCL_ASSERT(0);
> + }
> + }
> + OCL_UNMAP_BUFFER(1);
> + }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION (builtin_lgamma);
> diff --git a/utests/builtin_lgamma_r.cpp b/utests/builtin_lgamma_r.cpp
> new file mode 100644
> index 0000000..b6e5d0e
> --- /dev/null
> +++ b/utests/builtin_lgamma_r.cpp
> @@ -0,0 +1,46 @@
> +#include <cmath>
> +#include "utest_helper.hpp"
> +
> +void builtin_lgamma_r(void) {
> + const int n = 1024;
> + float src[n];
> +
> + // Setup kernel and buffers
> + OCL_CREATE_KERNEL("builtin_lgamma_r");
> + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
> + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
> + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int), NULL);
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
> + globals[0] = n;
> + locals[0] = 16;
> +
> + for (int j = 0; j < 1024; j++) {
> + OCL_MAP_BUFFER(0);
> + for (int i = 0; i < n; ++i) {
> + src[i] = ((float*) buf_data[0])[i] = (j * n + i + 1) * 0.001f;
> + }
> + OCL_UNMAP_BUFFER(0);
> +
> + OCL_NDRANGE(1);
> +
> + OCL_MAP_BUFFER(1);
> + OCL_MAP_BUFFER(2);
> + float *dst = (float*) buf_data[1];
> + for (int i = 0; i < n; ++i) {
> + int cpu_signp;
> + float cpu = lgamma_r(src[i], &cpu_signp);
> + int gpu_signp = ((int*)buf_data[2])[i];
> + float gpu = dst[i];
> + if (cpu_signp != gpu_signp || fabsf(cpu - gpu) >= 1e-3) {
> + printf("%f %f %f\n", src[i], cpu, gpu);
> + OCL_ASSERT(0);
> + }
> + }
> + OCL_UNMAP_BUFFER(1);
> + OCL_UNMAP_BUFFER(2);
> + }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION (builtin_lgamma_r);
> --
> 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