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

[Simon Richter]

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