diff options
author | Benjamin Segovia <benjamin.segovia@intel.com> | 2012-11-02 17:55:20 -0700 |
---|---|---|
committer | Benjamin Segovia <benjamin.segovia@intel.com> | 2012-11-02 17:55:20 -0700 |
commit | 4c5a0cfb128318adfdbc6d87cf77689584ee5bfb (patch) | |
tree | 7d68c78d5564a94abf145a86764a9413612a93d3 /kernels/compiler_clod.cl | |
parent | e4e590258d0d226fbf3193ef8a44c21379a8f61c (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.cl | 35 |
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 |