summaryrefslogtreecommitdiff
path: root/kernels/compiler_clod.cl
diff options
context:
space:
mode:
authorBenjamin Segovia <benjamin.segovia@intel.com>2012-11-02 17:55:20 -0700
committerBenjamin Segovia <benjamin.segovia@intel.com>2012-11-02 17:55:20 -0700
commit4c5a0cfb128318adfdbc6d87cf77689584ee5bfb (patch)
tree7d68c78d5564a94abf145a86764a9413612a93d3 /kernels/compiler_clod.cl
parente4e590258d0d226fbf3193ef8a44c21379a8f61c (diff)
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"
Diffstat (limited to 'kernels/compiler_clod.cl')
-rw-r--r--kernels/compiler_clod.cl35
1 files changed, 27 insertions, 8 deletions
diff --git a/kernels/compiler_clod.cl b/kernels/compiler_clod.cl
index d7c945a..ec9d33e 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