[Beignet] [PATCH 1/4] Define clamp(value, lower, upper)

[Zhigang Gong]

LGTM. I guess the reason that why the previous kernel was using a private clamp
is that when it was created, there is no clamp definition in ocl_stdlib.h:).

But it's still a little weird that the private definition always implicitly
use integer arguments. Anyway, this patch looks good for me. Will push it
soon. Thanks.

On Mon, May 13, 2013 at 08:21:15PM +0200, Simon Richter wrote:
> The clamp(value, lower, upper) function is part of the standard library.
> 
>  - Define the function, using min() and max() on the lower level
>  - Remove private definitions from kernels
> ---
>  backend/src/ocl_stdlib.h                    |   21 ++++++++++++---------
>  kernels/compiler_julia.cl                   |    2 --
>  kernels/compiler_julia_no_break.cl          |    2 --
>  kernels/compiler_menger_sponge.cl           |    2 --
>  kernels/compiler_menger_sponge_no_shadow.cl |    2 --
>  kernels/compiler_nautilus.cl                |    4 +---
>  6 files changed, 13 insertions(+), 20 deletions(-)
> 
> diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
> index 4c0d39c..7027373 100644
> --- a/backend/src/ocl_stdlib.h
> +++ b/backend/src/ocl_stdlib.h
> @@ -484,21 +484,24 @@ DECL_SELECT4(float4, float, uint4, 0x80000000)
>  /////////////////////////////////////////////////////////////////////////////
>  // Common Functions (see 6.11.4 of OCL 1.1 spec)
>  /////////////////////////////////////////////////////////////////////////////
> -#define DECL_MIN_MAX(TYPE) \
> +#define DECL_MIN_MAX_CLAMP(TYPE) \
>  INLINE OVERLOADABLE TYPE max(TYPE a, TYPE b) { \
>    return a > b ? a : b; \
>  } \
>  INLINE OVERLOADABLE TYPE min(TYPE a, TYPE b) { \
>    return a < b ? a : b; \
> +} \
> +INLINE OVERLOADABLE TYPE clamp(TYPE v, TYPE l, TYPE u) { \
> +  return max(min(v, u), l); \
>  }
> -DECL_MIN_MAX(float)
> -DECL_MIN_MAX(int)
> -DECL_MIN_MAX(short)
> -DECL_MIN_MAX(char)
> -DECL_MIN_MAX(uint)
> -DECL_MIN_MAX(unsigned short)
> -DECL_MIN_MAX(unsigned char)
> -#undef DECL_MIN_MAX
> +DECL_MIN_MAX_CLAMP(float)
> +DECL_MIN_MAX_CLAMP(int)
> +DECL_MIN_MAX_CLAMP(short)
> +DECL_MIN_MAX_CLAMP(char)
> +DECL_MIN_MAX_CLAMP(uint)
> +DECL_MIN_MAX_CLAMP(unsigned short)
> +DECL_MIN_MAX_CLAMP(unsigned char)
> +#undef DECL_MIN_MAX_CLAMP
>  
>  INLINE OVERLOADABLE float __gen_ocl_internal_fmax(float a, float b) { return max(a,b); }
>  INLINE OVERLOADABLE float __gen_ocl_internal_fmin(float a, float b) { return min(a,b); }
> diff --git a/kernels/compiler_julia.cl b/kernels/compiler_julia.cl
> index 996c0b7..98c5799 100644
> --- a/kernels/compiler_julia.cl
> +++ b/kernels/compiler_julia.cl
> @@ -14,8 +14,6 @@ inline vec3 reflect(vec3 I, vec3 N) {
>    return I - 2.0f * dot(N, I) * N;
>  }
>  
> -inline float clamp(x,m,M) { return max(min(x,M),m); }
> -
>  inline uint pack_fp4(float4 u4) {
>    uint u;
>    u = (((uint) u4.x)) |
> diff --git a/kernels/compiler_julia_no_break.cl b/kernels/compiler_julia_no_break.cl
> index c0bd3b1..1a9be64 100644
> --- a/kernels/compiler_julia_no_break.cl
> +++ b/kernels/compiler_julia_no_break.cl
> @@ -14,8 +14,6 @@ inline vec3 reflect(vec3 I, vec3 N) {
>    return I - 2.0f * dot(N, I) * N;
>  }
>  
> -inline float clamp(x,m,M) { return max(min(x,M),m); }
> -
>  inline uint pack_fp4(float4 u4) {
>    uint u;
>    u = (((uint) u4.x)) |
> diff --git a/kernels/compiler_menger_sponge.cl b/kernels/compiler_menger_sponge.cl
> index b59c5e3..58af12a 100644
> --- a/kernels/compiler_menger_sponge.cl
> +++ b/kernels/compiler_menger_sponge.cl
> @@ -25,8 +25,6 @@ inline vec3 reflect(vec3 I, vec3 N) {
>    return I - 2.0f * dot(N, I) * N;
>  }
>  
> -inline float clamp(x,m,M) { return max(min(x,M),m); }
> -
>  inline uint pack_fp4(float4 u4) {
>    uint u;
>    u = (((uint) u4.x)) |
> diff --git a/kernels/compiler_menger_sponge_no_shadow.cl b/kernels/compiler_menger_sponge_no_shadow.cl
> index 4f1093f..95469c5 100644
> --- a/kernels/compiler_menger_sponge_no_shadow.cl
> +++ b/kernels/compiler_menger_sponge_no_shadow.cl
> @@ -25,8 +25,6 @@ inline vec3 reflect(vec3 I, vec3 N) {
>    return I - 2.0f * dot(N, I) * N;
>  }
>  
> -inline float clamp(x,m,M) { return max(min(x,M),m); }
> -
>  inline uint pack_fp4(float4 u4) {
>    uint u;
>    u = (((uint) u4.x)) |
> diff --git a/kernels/compiler_nautilus.cl b/kernels/compiler_nautilus.cl
> index b53771c..aa7251a 100644
> --- a/kernels/compiler_nautilus.cl
> +++ b/kernels/compiler_nautilus.cl
> @@ -14,8 +14,6 @@ inline vec3 reflect(vec3 I, vec3 N) {
>    return I - 2.0f * dot(N, I) * N;
>  }
>  
> -inline float clamp(x,m,M) { return max(min(x,M),m); }
> -
>  inline uint pack_fp4(float4 u4) {
>    uint u;
>    u = (((uint) u4.x)) |
> @@ -59,7 +57,7 @@ __kernel void compiler_nautilus(__global uint *dst, float resx, float resy, int
>    for(int q=0;q<100;q++)
>    {
>       float l = e(o+0.5f*(vec3)(cos(1.1f*(float)(q)),cos(1.6f*(float)(q)),cos(1.4f*(float)(q))))-m;
> -     a+=clamp(4.0f*l,0.0f,1.0f);
> +     a+=floor(clamp(4.0f*l,0.0f,1.0f));
>    }
>    v*=a/100.0f;
>    vec4 gl_FragColor=(vec4)(v,1.0f);
> -- 
> 1.7.10.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet