Made compiler_clod pass. The image is now properly computed.

I basically added a bunch of new intrinsics and cleaned a bit the ocl std
library. Well, it is not going to be compliant for a while with the spec (mostly
due to precision issues and the way denormals and nan are handled). But,
it should do the job for now.
Started to add a more complex test called "ribbon"

2 files changed, 114 insertions, 8 deletions

diff --git a/kernels/compiler_clod.cl b/kernels/compiler_clod.cl
index d7c945ab..ec9d33e2 100644
--- a/kernels/compiler_clod.cl
+++ b/kernels/compiler_clod.cl
@@ -31,12 +31,13 @@ inline uint pack_fp4(float4 u4) {
 
 float f(vec3 o)
 {
-    float a=(sin(o.x)+o.y*.25)*.35f;
+    float a=(sin(o.x)+o.y*.25f)*.35f;
     o=(vec3)(cos(a)*o.x-sin(a)*o.y,sin(a)*o.x+cos(a)*o.y,o.z);
     return dot(cos(o)*cos(o),(vec3)(1.f))-1.2f;
 }
 
-vec3 s(vec3 o,vec3 d)
+// XXX front end does not inline this function
+__attribute((always_inline)) vec3 s(vec3 o,vec3 d)
 {
     float t=0.0f;
     float dt = 0.2f;
@@ -45,27 +46,45 @@ vec3 s(vec3 o,vec3 d)
     for(int i=0;i<50;i++)
     {
         nh = f(o+d*t);
-        if(nh>0.0) { lh=nh; t+=dt; }
+        if(nh>0.0f) { lh=nh; t+=dt; }
     }
 
-    if( nh>0.0 ) return (vec3)(.93f,.94f,.85f);
+    if( nh>0.0f ) return (vec3)(.93f,.94f,.85f);
 
     t = t - dt*nh/(nh-lh);
 
     vec3 e=(vec3)(.1f,0.0f,0.0f);
     vec3 p=o+d*t;
-    vec3 n=-normalize((vec3)(f(p+e),f(p+e.yxy),f(p+e.yyx))+(vec3)((sin(p*75.)))*.01f);
+    vec3 n=-normalize((vec3)(f(p+e),f(p+e.yxy),f(p+e.yyx))+(vec3)((sin(p*75.f)))*.01f);
 
     return (vec3)(mix( ((max(-dot(n,(vec3)(.577f)),0.f) + 0.125f*max(-dot(n,(vec3)(-.707f,-.707f,0.f)),0.f)))*(mod
-    (length(p.xy)*20.f,2.f)<1.0?(vec3)(.71f,.85f,.25f):(vec3)(.79f,.93f,.4f))
+    (length(p.xy)*20.f,2.f)<1.0f?(vec3)(.71f,.85f,.25f):(vec3)(.79f,.93f,.4f))
                            ,(vec3)(.93f,.94f,.85f), (vec3)(pow(t/9.f,5.f)) ) );
 }
 
+#if 0
+// XXX vector type in the function arguments not supported yet
 __kernel void compiler_clod(__global uint *dst, vec2 resolution, int w)
 {
     vec2 gl_FragCoord = (vec2)(get_global_id(0), get_global_id(1));
-    vec2 p = -1.0f + 2.0f * gl_FragCoord.xy / resolution.xy;
-    vec4 gl_FragColor=(vec4)(s((vec3)(sin(time*1.5)*.5f,cos(time)*.5f,time), normalize((vec3)(p.xy,1.0f))),1.0f);
+    //vec2 p = -1.0f + 2.0f * gl_FragCoord.xy / resolution.xy;
+    vec2 p;
+    p.x = -1.0f + 2.0f * gl_FragCoord.x / resolution.x;
+    p.y = -1.0f + 2.0f * gl_FragCoord.y / resolution.y;
+    vec4 gl_FragColor=(vec4)(s((vec3)(sin(time*1.5f)*.5f,cos(time)*.5f,time), normalize((vec3)(p.xy,1.0f))),1.0f);
     OUTPUT;
 }
+#else
+__kernel void compiler_clod(__global uint *dst, float resx, float resy, int w)
+{
+    vec2 gl_FragCoord = (vec2)(get_global_id(0), get_global_id(1));
+    //vec2 p = -1.0f + 2.0f * gl_FragCoord.xy / resolution.xy;
+    vec2 p;
+    p.x = -1.0f + 2.0f * gl_FragCoord.x / resx;
+    p.y = -1.0f + 2.0f * gl_FragCoord.y / resy;
+    vec4 gl_FragColor=(vec4)(s((vec3)(sin(time*1.5f)*.5f,cos(time)*.5f,time), normalize((vec3)(p.xy,1.0f))),1.0f);
+    OUTPUT;
+}
+
+#endif
 
diff --git a/kernels/compiler_ribbon.cl b/kernels/compiler_ribbon.cl
new file mode 100644
index 00000000..f5729e55
--- /dev/null
+++ b/kernels/compiler_ribbon.cl
@@ -0,0 +1,87 @@
+typedef float2 vec2;

+typedef float3 vec3;

+typedef float4 vec4;

+#define sin native_sin

+#define cos native_cos

+#define tan native_tan

+#define normalize fast_normalize

+#define length fast_length

+

+inline vec3 reflect(vec3 I, vec3 N) {

+  return I - 2.0f * dot(N, I) * N;

+}

+

+#define time 1.f

+

+// Object A (tunnel)

+inline float oa(vec3 q) {

+ return cos(q.x)+cos(q.y*1.5f)+cos(q.z)+cos(q.y*20.f)*.05f;

+}

+

+// Object B (ribbon)

+inline float ob(vec3 q) {

+  return length(max(fabs(q-(vec3)(cos(q.z*1.5f)*.3f,-.5f+cos(q.z)*.2f,.0f))-(vec3)(.125f,.02f,time+3.f),(vec3)(.0f)));

+}

+

+// Scene

+inline float o(vec3 q) { return min(oa(q),ob(q)); }

+

+// Get Normal XXX Not inline by LLVM

+__attribute__((always_inline)) vec3 gn(vec3 q) {

+ const vec3 f = (vec3)(.01f, 0.f, 0.f);

+ return normalize((vec3)(o(q+f.xyy),

+                         o(q+f.yxy),

+                         o(q+f.yyx)));

+}

+

+inline uint pack_fp4(float4 u4) {

+  uint u;

+  u = (((uint) u4.x)) |

+      (((uint) u4.y) << 8) |

+      (((uint) u4.z) << 16);

+  return u;

+}

+

+// XXX vector not supported in function argument yet

+__kernel void compiler_ribbon(__global uint *dst, float resx, float resy, int w)

+{

+  vec2 gl_FragCoord = (vec2)(get_global_id(0), get_global_id(1));

+  vec2 p = -1.0f + 2.0f * gl_FragCoord.xy / (vec2)(resx, resy);

+  p.x *= resx/resy;

+

+  vec4 c = (vec4)(1.0f);

+  const vec3 org = (vec3)(sin(time)*.5f,

+                          cos(time*.5f)*.25f+.25f,

+                          time);

+  vec3 dir=normalize((vec3)(p.x*1.6f,p.y,1.0f));

+  vec3 q = org, pp;

+  float d=.0f;

+

+  // First raymarching

+  for(int i=0;i<64;i++) {

+    d=o(q);

+    q+=d*dir;

+  }

+  pp=q;

+  const float f = length(q-org)*0.02f;

+

+  // Second raymarching (reflection)

+  dir=reflect(dir,gn(q));

+  q+=dir;

+  for(int i=0;i<64;i++) {

+    d=o(q);

+    q+=d*dir;

+  }

+  c = max(dot(gn(q), (vec3)(0.1f,0.1f,0.0f)), 0.0f)

+    + (vec4)(0.3f, cos(time*.5f)*.5f+.5f, sin(time*.5f)*.5f+.5f, 1.f) * min(length(q-org)*.04f,1.f);

+

+  // Ribbon Color

+  if(oa(pp)>ob(pp))

+    c = mix(c, (vec4)(cos(time*.3f)*0.5f + 0.5f,cos(time*.2f)*.5f+.5f,sin(time*.3f)*.5f+.5f,1.f),.3f);

+

+  // Final Color

+  const vec4 color = ((c+(vec4)(f))+(1.f-min(pp.y+1.9f,1.f))*(vec4)(1.f,.8f,.7f,1.f))*min(time*.5f,1.f);

+  const vec4 final = 255.f * max(min(color, (vec4)(1.f)), (vec4)(0.f));

+  dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final);

+}

+