diff options
-rw-r--r-- | kernels/compiler_chocolux.cl | 64 | ||||
-rw-r--r-- | kernels/compiler_chocolux_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_clod.cl | 91 | ||||
-rw-r--r-- | kernels/compiler_clod_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_julia.cl | 144 | ||||
-rw-r--r-- | kernels/compiler_julia_function_call.cl | 142 | ||||
-rw-r--r-- | kernels/compiler_julia_no_break.cl | 145 | ||||
-rw-r--r-- | kernels/compiler_julia_no_break_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_julia_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_menger_sponge.cl | 187 | ||||
-rw-r--r-- | kernels/compiler_menger_sponge_no_shadow.cl | 123 | ||||
-rw-r--r-- | kernels/compiler_menger_sponge_no_shadow_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_menger_sponge_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_nautilus.cl | 66 | ||||
-rw-r--r-- | kernels/compiler_nautilus_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | kernels/compiler_ribbon.cl | 88 | ||||
-rw-r--r-- | kernels/compiler_ribbon_ref.bmp | bin | 196662 -> 0 bytes | |||
-rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
-rw-r--r-- | utests/compiler_shader_toy.cpp | 87 |
19 files changed, 0 insertions, 1138 deletions
diff --git a/kernels/compiler_chocolux.cl b/kernels/compiler_chocolux.cl deleted file mode 100644 index 218f65db..00000000 --- a/kernels/compiler_chocolux.cl +++ /dev/null @@ -1,64 +0,0 @@ -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 -#define mod fmod -#define time 10.f - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -__kernel void compiler_chocolux(__global uint *dst, float resx, float resy, int w) -{ - vec2 gl_FragCoord = (vec2)(get_global_id(0), get_global_id(1)); - vec3 s[4]; - s[0]=(vec3)(0); - s[3]=(vec3)(sin(time),cos(time),0); - s[1]=s[3].zxy; - s[2]=s[3].zzx; - - float t,b,c,h=0.0f; - vec3 m,n; - vec3 p=(vec3)(.2f); - vec3 d=normalize(.001f*(vec3)(gl_FragCoord,.0f)-p); - - for(int i=0;i<4;i++) - { - t=2.0f; - for(int i=0;i<4;i++) - { - b=dot(d,n=s[i]-p); - c=b*b+.2f-dot(n,n); - if(b-c<t) - if(c>0.0f) - { - m=s[i];t=b-c; - } - } - p+=t*d; - d=reflect(d,n=normalize(p-m)); - h+=pow(n.x*n.x,44.f)+n.x*n.x*.2f; - } - vec4 gl_FragColor=(vec4)(h,h*h,h*h*h*h,1.f); - OUTPUT; -} - diff --git a/kernels/compiler_chocolux_ref.bmp b/kernels/compiler_chocolux_ref.bmp Binary files differdeleted file mode 100644 index e51a4a73..00000000 --- a/kernels/compiler_chocolux_ref.bmp +++ /dev/null diff --git a/kernels/compiler_clod.cl b/kernels/compiler_clod.cl deleted file mode 100644 index dba7d6fb..00000000 --- a/kernels/compiler_clod.cl +++ /dev/null @@ -1,91 +0,0 @@ -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 -#define mod fmod - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -#define time 1.f - -inline float f(vec3 o) -{ - 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; -} - -// XXX front end does not inline this function -inline __attribute((always_inline)) vec3 s(vec3 o,vec3 d) -{ - float t=0.0f; - float dt = 0.2f; - float nh = 0.0f; - float lh = 0.0f; - for(int i=0;i<50;i++) - { - nh = f(o+d*t); - if(nh>0.0f) { lh=nh; t+=dt; } - } - - if( nh>0.0f ) return (vec3)(.93f,.94f,.85f); - - t = t - dt*nh/(nh-lh); - - vec3 exyy=(vec3)(0.1f,0.0f,0.0f); - vec3 eyxy=(vec3)(0.0f,0.1f,0.0f); - vec3 eyyx=(vec3)(0.0f,0.0f,0.1f); - vec3 p=o+d*t; - vec3 n=-normalize((vec3)(f(p+exyy),f(p+eyxy),f(p+eyyx))+(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.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; - 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_clod_ref.bmp b/kernels/compiler_clod_ref.bmp Binary files differdeleted file mode 100644 index 71afda90..00000000 --- a/kernels/compiler_clod_ref.bmp +++ /dev/null diff --git a/kernels/compiler_julia.cl b/kernels/compiler_julia.cl deleted file mode 100644 index 21672f69..00000000 --- a/kernels/compiler_julia.cl +++ /dev/null @@ -1,144 +0,0 @@ -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 -#define mod fmod -#define time 1.f - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -inline __attribute__((always_inline)) -float jinteresct(vec3 rO, vec3 rD, vec4 c, float *ao) -{ - float mz2,md2,dist,t; - float res=1000.0f; - vec4 z,nz; - int update_ao = 1; - *ao = 0.0f; - for(t=0.0f;t<6.0f;t+=dist) - { - if (update_ao) *ao += 1.0f; - vec3 p=rO+t*rD; - - // calc distance - z=(vec4)(p,(c.y+c.x)*.3f); - md2=1.0f; - mz2=dot(z,z); - - for(int i=0;i<9;i++) - { - // |dz|^2 -> 4*|dz|^2 - //if (mz2 <= 4.0f) - { - md2*=4.0f*mz2; - // z -> z2 + c - nz.x=z.x*z.x-dot(z.yzw,z.yzw); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - mz2=dot(z,z); - } - if(mz2>4.0f) - break; - } - - dist=0.25f*sqrt(mz2/md2)*log(mz2); - if(dist<0.0005f) - { - res=t; - break; - } - t+= dist; - } - - return res; -} - -#if 1 -inline __attribute__((always_inline)) -vec3 calcNormal(vec3 p, vec4 c) -{ - vec4 nz,ndz,dz[4]; - - vec4 z=(vec4)(p,(c.y+c.x)*.3f); - - dz[0]=(vec4)(1.0f,0.0f,0.0f,0.0f); - dz[1]=(vec4)(0.0f,1.0f,0.0f,0.0f); - dz[2]=(vec4)(0.0f,0.0f,1.0f,0.0f); - //dz[3]=(vec4)(0.0f,0.0f,0.0f,1.0f); - - for(int i=0;i<9;i++) - { - vec4 mz = (vec4)(z.x,-z.y,-z.z,-z.w); - // derivative - dz[0]=(vec4)(dot(mz,dz[0]),z.x*dz[0].yzw+dz[0].x*z.yzw); - dz[1]=(vec4)(dot(mz,dz[1]),z.x*dz[1].yzw+dz[1].x*z.yzw); - dz[2]=(vec4)(dot(mz,dz[2]),z.x*dz[2].yzw+dz[2].x*z.yzw); - //dz[3]=(vec4)(dot(mz,dz[3]),z.x*dz[3].yzw+dz[3].x*z.yzw); - - // z = z2 + c - nz.x=dot(z, mz); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - - if(dot(z,z)>4.0f) - break; - } - - return normalize((vec3)(dot(z,dz[0]),dot(z,dz[1]),dot(z,dz[2]))); -} -#endif - -__kernel void compiler_julia(__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); - vec3 color = (vec3)(0.0f); - vec4 cccc = (vec4)( .7f*cos(.5f*time), .7f*sin(.3f*time), .7f*cos(1.0f*time), 0.0f ); - vec3 edir = normalize((vec3)(p,1.0f)); - vec3 wori = (vec3)(0.0f,0.0f,-2.0f); - - float ao; - float t = jinteresct(wori,edir,cccc,&ao); - if(t<100.0f) - { -#if 1 - vec3 inter = wori + t*edir; - vec3 nor = calcNormal(inter,cccc); - - float dif = .5f + .5f*dot( nor, (vec3)(0.57703f) ); - ao = max( 1.0f-ao*0.005f, 0.0f); - - color = (vec3)(1.0f,.9f,.5f)*dif*ao + .5f*(vec3)(.6f,.7f,.8f)*ao; -#else - color = (vec3)(0.5f,0.0f,0.0f); -#endif - } - else - { - color = (vec3)(0.5f,0.51f,0.52f)+(vec3)(0.5f,0.47f,0.45f)*p.y; - } - - vec4 gl_FragColor = (vec4)(color,1.0f); - OUTPUT; -} - diff --git a/kernels/compiler_julia_function_call.cl b/kernels/compiler_julia_function_call.cl deleted file mode 100644 index 7b3aa464..00000000 --- a/kernels/compiler_julia_function_call.cl +++ /dev/null @@ -1,142 +0,0 @@ -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 -#define mod fmod -#define time 1.f - -vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -float jinteresct(vec3 rO, vec3 rD, vec4 c, float *ao) -{ - float mz2,md2,dist,t; - float res=1000.0f; - vec4 z,nz; - int update_ao = 1; - *ao = 0.0f; - for(t=0.0f;t<6.0f;t+=dist) - { - if (update_ao) *ao += 1.0f; - vec3 p=rO+t*rD; - - // calc distance - z=(vec4)(p,(c.y+c.x)*.3f); - md2=1.0f; - mz2=dot(z,z); - - for(int i=0;i<9;i++) - { - // |dz|^2 -> 4*|dz|^2 - //if (mz2 <= 4.0f) - { - md2*=4.0f*mz2; - // z -> z2 + c - nz.x=z.x*z.x-dot(z.yzw,z.yzw); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - mz2=dot(z,z); - } - if(mz2>4.0f) - break; - } - - dist=0.25f*sqrt(mz2/md2)*log(mz2); - if(dist<0.0005f) - { - res=t; - break; - } - t+= dist; - } - - return res; -} - -#if 1 -vec3 calcNormal(vec3 p, vec4 c) -{ - vec4 nz,ndz,dz[4]; - - vec4 z=(vec4)(p,(c.y+c.x)*.3f); - - dz[0]=(vec4)(1.0f,0.0f,0.0f,0.0f); - dz[1]=(vec4)(0.0f,1.0f,0.0f,0.0f); - dz[2]=(vec4)(0.0f,0.0f,1.0f,0.0f); - //dz[3]=(vec4)(0.0f,0.0f,0.0f,1.0f); - - for(int i=0;i<9;i++) - { - vec4 mz = (vec4)(z.x,-z.y,-z.z,-z.w); - // derivative - dz[0]=(vec4)(dot(mz,dz[0]),z.x*dz[0].yzw+dz[0].x*z.yzw); - dz[1]=(vec4)(dot(mz,dz[1]),z.x*dz[1].yzw+dz[1].x*z.yzw); - dz[2]=(vec4)(dot(mz,dz[2]),z.x*dz[2].yzw+dz[2].x*z.yzw); - //dz[3]=(vec4)(dot(mz,dz[3]),z.x*dz[3].yzw+dz[3].x*z.yzw); - - // z = z2 + c - nz.x=dot(z, mz); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - - if(dot(z,z)>4.0f) - break; - } - - return normalize((vec3)(dot(z,dz[0]),dot(z,dz[1]),dot(z,dz[2]))); -} -#endif - -__kernel void compiler_julia(__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); - vec3 color = (vec3)(0.0f); - vec4 cccc = (vec4)( .7f*cos(.5f*time), .7f*sin(.3f*time), .7f*cos(1.0f*time), 0.0f ); - vec3 edir = normalize((vec3)(p,1.0f)); - vec3 wori = (vec3)(0.0f,0.0f,-2.0f); - - float ao; - float t = jinteresct(wori,edir,cccc,&ao); - if(t<100.0f) - { -#if 1 - vec3 inter = wori + t*edir; - vec3 nor = calcNormal(inter,cccc); - - float dif = .5f + .5f*dot( nor, (vec3)(0.57703f) ); - ao = max( 1.0f-ao*0.005f, 0.0f); - - color = (vec3)(1.0f,.9f,.5f)*dif*ao + .5f*(vec3)(.6f,.7f,.8f)*ao; -#else - color = (vec3)(0.5f,0.0f,0.0f); -#endif - } - else - { - color = (vec3)(0.5f,0.51f,0.52f)+(vec3)(0.5f,0.47f,0.45f)*p.y; - } - - vec4 gl_FragColor = (vec4)(color,1.0f); - OUTPUT; -} - diff --git a/kernels/compiler_julia_no_break.cl b/kernels/compiler_julia_no_break.cl deleted file mode 100644 index 5c357b1d..00000000 --- a/kernels/compiler_julia_no_break.cl +++ /dev/null @@ -1,145 +0,0 @@ -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 -#define mod fmod -#define time 1.f - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -inline __attribute__((always_inline)) -float jinteresct(vec3 rO, vec3 rD, vec4 c, float *ao) -{ - float mz2,md2,dist,t; - float res=1000.0f; - vec4 z,nz; - int update_ao = 1; - *ao = 0.0f; - t = 0.f; - for (int j = 0; j < 100; ++j) - { - if (update_ao) *ao += 1.0f; - vec3 p=rO+t*rD; - - // calc distance - z=(vec4)(p,(c.y+c.x)*.3f); - md2=1.0f; - mz2=dot(z,z); - - for(int i=0;i<9;i++) - { - // |dz|^2 -> 4*|dz|^2 - //if (mz2 <= 4.0f) - { - md2*=4.0f*mz2; - // z -> z2 + c - nz.x=z.x*z.x-dot(z.yzw,z.yzw); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - mz2=dot(z,z); - } - if(mz2>4.0f) - break; - } - - dist=0.25f*sqrt(mz2/md2)*log(mz2); - if(dist<0.0005f) - { - res=t; - update_ao = 0; - } - t+= dist; - } - - return res; -} - -#if 1 -inline __attribute__((always_inline)) -vec3 calcNormal(vec3 p, vec4 c) -{ - vec4 nz,ndz,dz[4]; - - vec4 z=(vec4)(p,(c.y+c.x)*.3f); - - dz[0]=(vec4)(1.0f,0.0f,0.0f,0.0f); - dz[1]=(vec4)(0.0f,1.0f,0.0f,0.0f); - dz[2]=(vec4)(0.0f,0.0f,1.0f,0.0f); - //dz[3]=(vec4)(0.0f,0.0f,0.0f,1.0f); - - for(int i=0;i<9;i++) - { - vec4 mz = (vec4)(z.x,-z.y,-z.z,-z.w); - // derivative - dz[0]=(vec4)(dot(mz,dz[0]),z.x*dz[0].yzw+dz[0].x*z.yzw); - dz[1]=(vec4)(dot(mz,dz[1]),z.x*dz[1].yzw+dz[1].x*z.yzw); - dz[2]=(vec4)(dot(mz,dz[2]),z.x*dz[2].yzw+dz[2].x*z.yzw); - //dz[3]=(vec4)(dot(mz,dz[3]),z.x*dz[3].yzw+dz[3].x*z.yzw); - - // z = z2 + c - nz.x=dot(z, mz); - nz.yzw=2.0f*z.x*z.yzw; - z=nz+c; - - if(dot(z,z)>4.0f) - break; - } - - return normalize((vec3)(dot(z,dz[0]),dot(z,dz[1]),dot(z,dz[2]))); -} -#endif - -__kernel void compiler_julia_no_break(__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); - vec3 color = (vec3)(0.0f); - vec4 cccc = (vec4)( .7f*cos(.5f*time), .7f*sin(.3f*time), .7f*cos(1.0f*time), 0.0f ); - vec3 edir = normalize((vec3)(p,1.0f)); - vec3 wori = (vec3)(0.0f,0.0f,-2.0f); - - float ao; - float t = jinteresct(wori,edir,cccc,&ao); - if(t<100.0f) - { -#if 1 - vec3 inter = wori + t*edir; - vec3 nor = calcNormal(inter,cccc); - - float dif = .5f + .5f*dot( nor, (vec3)(0.57703f) ); - ao = max( 1.0f-ao*0.005f, 0.0f); - - color = (vec3)(1.0f,.9f,.5f)*dif*ao + .5f*(vec3)(.6f,.7f,.8f)*ao; -#else - color = (vec3)(0.5f,0.0f,0.0f); -#endif - } - else - { - color = (vec3)(0.5f,0.51f,0.52f)+(vec3)(0.5f,0.47f,0.45f)*p.y; - } - - vec4 gl_FragColor = (vec4)(color,1.0f); - OUTPUT; -} - diff --git a/kernels/compiler_julia_no_break_ref.bmp b/kernels/compiler_julia_no_break_ref.bmp Binary files differdeleted file mode 100644 index e17f6660..00000000 --- a/kernels/compiler_julia_no_break_ref.bmp +++ /dev/null diff --git a/kernels/compiler_julia_ref.bmp b/kernels/compiler_julia_ref.bmp Binary files differdeleted file mode 100644 index 2082a1e9..00000000 --- a/kernels/compiler_julia_ref.bmp +++ /dev/null diff --git a/kernels/compiler_menger_sponge.cl b/kernels/compiler_menger_sponge.cl deleted file mode 100644 index 58af12a8..00000000 --- a/kernels/compiler_menger_sponge.cl +++ /dev/null @@ -1,187 +0,0 @@ -// See http://www.iquilezles.org/articles/menger/menger.htm for the -// full explanation of how this was done - -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 -#define mod fmod -#define time 1.f - -// fmod is not like glsl mod! -__attribute__((always_inline, overloadable)) -float glsl_mod(float x,float y) { return x-y*floor(x/y); } -__attribute__((always_inline, overloadable)) -float2 glsl_mod(float2 a,float2 b) { return (float2)(glsl_mod(a.x,b.x), glsl_mod(a.y,b.y)); } -__attribute__((always_inline, overloadable)) -float3 glsl_mod(float3 a,float3 b) { return (float3)(glsl_mod(a.x,b.x), glsl_mod(a.y,b.y), glsl_mod(a.z,b.z)); } - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -__attribute__((always_inline)) -float maxcomp(vec3 p) { return max(p.x,max(p.y,p.z));} - -__attribute__((always_inline)) -float sdBox(vec3 p, vec3 b) -{ - vec3 di = fabs(p) - b; - float mc = maxcomp(di); - return min(mc,length(max(di,0.0f))); -} - -__attribute__((always_inline)) -vec4 map(vec3 p) -{ - float d = sdBox(p,(vec3)(1.0f)); - float4 res = (vec4)(d,1.f,0.f,0.f); - - float s = 1.0f; - for( int m=0; m<3; m++ ) - { - vec3 a = glsl_mod(p*s, 2.0f)-1.0f; - s *= 3.0f; - float rx = fabs(1.0f - 3.0f*fabs(a.x)); - float ry = fabs(1.0f - 3.0f*fabs(a.y)); - float rz = fabs(1.0f - 3.0f*fabs(a.z)); - - float da = max(rx,ry); - float db = max(ry,rz); - float dc = max(rz,rx); - float c = (min(da,min(db,dc))-1.0f)/s; - if (c > d) - { - d = c; - res = (vec4)(d, 0.2f*da*db*dc, (1.0f+(float)(m))/4.0f, 0.0f); - } - } - return (vec4)(res.x,res.y,res.z,0.f); -} - -// GLSL ES doesn't seem to like loops with conditional break/return... -#if 1 -__attribute__((always_inline)) -vec4 intersect( vec3 ro, vec3 rd ) -{ - float t = 0.0f; - for(int i=0;i<64;i++) - { - vec4 h = map(ro + rd*t); - if( h.x<0.002f ) - return (vec4)(t,h.yzw); - t += h.x; - } - return (vec4)(-1.0f); -} -#else -__attribute__((always_inline)) -vec4 intersect( vec3 ro, vec3 rd ) -{ - float t = 0.0f; - vec4 res = (vec4)(-1.0f); - for(int i=0;i<64;i++) - { - vec4 h = map(ro + rd*t); - if (h.x<0.002f) - { - if(res.x<0.0f) res = (vec4)(t,h.yzw); - } - t += h.x; - } - return res; -} -#endif - -__attribute__((always_inline)) -vec3 calcNormal(vec3 pos) -{ - vec3 epsxyy = (vec3)(.001f,0.0f,0.0f); - vec3 epsyxy = (vec3)(0.0f,.001f,0.0f); - vec3 epsyyx = (vec3)(0.0f,0.0f,.001f); - vec3 nor; - nor.x = map(pos+epsxyy).x - map(pos-epsxyy).x; - nor.y = map(pos+epsyxy).x - map(pos-epsyxy).x; - nor.z = map(pos+epsyyx).x - map(pos-epsyyx).x; - return normalize(nor); -} - -__kernel void compiler_menger_sponge(__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); - - // light - vec3 light = normalize((vec3)(1.0f,0.8f,-0.6f)); - - float ctime = time; - // camera - vec3 ro = 1.1f*(vec3)(2.5f*cos(0.5f*ctime),1.5f*cos(ctime*.23f),2.5f*sin(0.5f*ctime)); - vec3 ww = normalize((vec3)(0.0f) - ro); - vec3 uu = normalize(cross( (vec3)(0.0f,1.0f,0.0f), ww )); - vec3 vv = normalize(cross(ww,uu)); - vec3 rd = normalize( p.x*uu + p.y*vv + 1.5f*ww ); - vec3 col = (vec3)(0.0f); - vec4 tmat = intersect(ro,rd); - -#if 0 - if( tmat.x>0.0 ) - col = (vec3)( - 0.6f+0.4f*cos(5.0f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.4f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.7f+6.2831f*tmat.z) ); - -#else - if( tmat.x>0.0f ) - { - vec3 pos = ro + tmat.x*rd; - vec3 nor = calcNormal(pos); - - float dif1 = max(0.4f + 0.6f*dot(nor,light),0.0f); - float dif2 = max(0.4f + 0.6f*dot(nor,(vec3)(-light.x,light.y,-light.z)),0.0f); - - // shadow - float ldis = 4.0f; - vec4 shadow = intersect( pos + light*ldis, -light ); - if( shadow.x>0.0f && shadow.x<(ldis-0.01f) ) dif1=0.0f; - - float ao = tmat.y; - col = 1.0f*ao*(vec3) (0.2f,0.2f,0.2f); - col += 2.0f*(0.5f+0.5f*ao)*dif1*(vec3)(1.0f,0.97f,0.85f); - col += 0.2f*(0.5f+0.5f*ao)*dif2*(vec3)(1.0f,0.97f,0.85f); - col += 1.0f*(0.5f+0.5f*ao)*(0.5f+0.5f*nor.y)*(vec3)(0.1f,0.15f,0.2f); - - // gamma lighting - col = col*0.5f+0.5f*sqrt(col)*1.2f; - - vec3 matcol = (vec3)( - 0.6f+0.4f*cos(5.0f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.4f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.7f+6.2831f*tmat.z) ); - col *= matcol; - col *= 1.5f*exp(-0.5f*tmat.x); - } -#endif - - vec4 gl_FragColor = (vec4)(col,1.0f); - OUTPUT; -} - diff --git a/kernels/compiler_menger_sponge_no_shadow.cl b/kernels/compiler_menger_sponge_no_shadow.cl deleted file mode 100644 index 27b059a5..00000000 --- a/kernels/compiler_menger_sponge_no_shadow.cl +++ /dev/null @@ -1,123 +0,0 @@ -// See http://www.iquilezles.org/articles/menger/menger.htm for the -// full explanation of how this was done - -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 -#define mod fmod -#define time 1.f - -// fmod is not like glsl mod! -inline __attribute__((always_inline, overloadable)) -float glsl_mod(float x,float y) { return mad( -y, floor(x/y), x); } -inline __attribute__((always_inline, overloadable)) -float2 glsl_mod(float2 a,float2 b) { return (float2)(glsl_mod(a.x,b.x), glsl_mod(a.y,b.y)); } -inline __attribute__((always_inline, overloadable)) -float3 glsl_mod(float3 a,float3 b) { return (float3)(glsl_mod(a.x,b.x), glsl_mod(a.y,b.y), glsl_mod(a.z,b.z)); } - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -inline __attribute__((always_inline)) -float maxcomp(vec3 p) { return max(p.x,max(p.y,p.z));} - -inline __attribute__((always_inline)) -float sdBox(vec3 p, vec3 b) -{ - vec3 di = fabs(p) - b; - float mc = maxcomp(di); - return min(mc,length(max(di,0.0f))); -} - -inline __attribute__((always_inline)) -vec4 map(vec3 p) -{ - float d = sdBox(p,(vec3)(1.0f)); - float4 res = (vec4)(d,1.f,0.f,0.f); - - float s = 1.0f; - for( int m=0; m<3; m++ ) - { - vec3 a = glsl_mod(p*s, 2.0f)-1.0f; - s *= 3.0f; - float rx = fabs(1.0f - 3.0f*fabs(a.x)); - float ry = fabs(1.0f - 3.0f*fabs(a.y)); - float rz = fabs(1.0f - 3.0f*fabs(a.z)); - - float da = max(rx,ry); - float db = max(ry,rz); - float dc = max(rz,rx); - float c = (min(da,min(db,dc))-1.0f)/s; - if (c > d) - { - d = c; - res = (vec4)(d, 0.2f*da*db*dc, (1.0f+(float)(m))/4.0f, 0.0f); - } - } - return (vec4)(res.x,res.y,res.z,0.f); -} - -// GLSL ES doesn't seem to like loops with conditional break/return... -inline __attribute__((always_inline)) -vec4 intersect( vec3 ro, vec3 rd ) -{ - float t = 0.0f; - for(int i=0;i<64;i++) - { - vec4 h = map(ro + rd*t); - if( h.x<0.002f ) - return (vec4)(t,h.yzw); - t += h.x; - } - return (vec4)(-1.0f); -} - -__kernel void compiler_menger_sponge_no_shadow(__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); - - // light - vec3 light = normalize((vec3)(1.0f,0.8f,-0.6f)); - - float ctime = time; - // camera - vec3 ro = 1.1f*(vec3)(2.5f*cos(0.5f*ctime),1.5f*cos(ctime*.23f),2.5f*sin(0.5f*ctime)); - vec3 ww = normalize((vec3)(0.0f) - ro); - vec3 uu = normalize(cross( (vec3)(0.0f,1.0f,0.0f), ww )); - vec3 vv = normalize(cross(ww,uu)); - vec3 rd = normalize( p.x*uu + p.y*vv + 1.5f*ww ); - vec3 col = (vec3)(0.0f); - vec4 tmat = intersect(ro,rd); - - if( tmat.x>0.0f ) - col = (vec3)( - 0.6f+0.4f*cos(5.0f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.4f+6.2831f*tmat.z), - 0.6f+0.4f*cos(5.7f+6.2831f*tmat.z) ); - - vec4 gl_FragColor = (vec4)(col,1.0f); - OUTPUT; -} - - diff --git a/kernels/compiler_menger_sponge_no_shadow_ref.bmp b/kernels/compiler_menger_sponge_no_shadow_ref.bmp Binary files differdeleted file mode 100644 index 133dd1d0..00000000 --- a/kernels/compiler_menger_sponge_no_shadow_ref.bmp +++ /dev/null diff --git a/kernels/compiler_menger_sponge_ref.bmp b/kernels/compiler_menger_sponge_ref.bmp Binary files differdeleted file mode 100644 index 911289fd..00000000 --- a/kernels/compiler_menger_sponge_ref.bmp +++ /dev/null diff --git a/kernels/compiler_nautilus.cl b/kernels/compiler_nautilus.cl deleted file mode 100644 index aa7251a8..00000000 --- a/kernels/compiler_nautilus.cl +++ /dev/null @@ -1,66 +0,0 @@ -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 -#define mod fmod -#define time 1.f - -inline vec3 reflect(vec3 I, vec3 N) { - return I - 2.0f * dot(N, I) * N; -} - -inline uint pack_fp4(float4 u4) { - uint u; - u = (((uint) u4.x)) | - (((uint) u4.y) << 8) | - (((uint) u4.z) << 16); - return u; -} - -#define OUTPUT do {\ - const vec4 final = 255.f * max(min(gl_FragColor, (vec4)(1.f)), (vec4)(0.f)); \ - dst[get_global_id(0) + get_global_id(1) * w] = pack_fp4(final); \ -} while (0) - -inline float e(vec3 c) -{ - c=cos((vec3)(cos(c.x+time/6.0f)*c.x-cos(c.y*3.0f+time/5.0f)*c.y, - cos(time/4.0f)*c.z/3.0f*c.x-cos(time/7.0f)*c.y, - c.x+c.y+c.z+time)); - return dot(c*c,(vec3)(1.0f))-1.0f; -} - -__kernel void compiler_nautilus(__global uint *dst, float resx, float resy, int w) -{ - vec2 gl_FragCoord = (vec2)(get_global_id(0), get_global_id(1)); - vec2 c=-1.0f+2.0f*gl_FragCoord.xy/(vec2)(resx,resy); - vec3 o=(vec3)(c.x,c.y,0.0f),g=(vec3)(c.x,c.y,1.0f)/64.0f,v=(vec3)(0.5f); - float m = 0.4f; - - for(int r=0;r<100;r++) - { - float h=e(o)-m; - if(h<0.0f)break; - o+=h*10.0f*g; - v+=h*0.02f; - } - // light (who needs a normal?) - v+=e(o+0.1f)*(vec3)(0.4f,0.7f,1.0f); - - // ambient occlusion - float a=0.0f; - 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+=floor(clamp(4.0f*l,0.0f,1.0f)); - } - v*=a/100.0f; - vec4 gl_FragColor=(vec4)(v,1.0f); - OUTPUT; -} - diff --git a/kernels/compiler_nautilus_ref.bmp b/kernels/compiler_nautilus_ref.bmp Binary files differdeleted file mode 100644 index 9d2dd96c..00000000 --- a/kernels/compiler_nautilus_ref.bmp +++ /dev/null diff --git a/kernels/compiler_ribbon.cl b/kernels/compiler_ribbon.cl deleted file mode 100644 index 157cc66a..00000000 --- a/kernels/compiler_ribbon.cl +++ /dev/null @@ -1,88 +0,0 @@ -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 -inline __attribute__((always_inline)) vec3 gn(vec3 q) { - const vec3 fxyy = (vec3)(.01f, 0.f, 0.f); - const vec3 fyxy = (vec3)(0.f, .01f, 0.f); - const vec3 fyyx = (vec3)(0.f, 0.f, .01f); - return normalize((vec3)(o(q+fxyy), - o(q+fyxy), - o(q+fyyx))); -} - -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); -} diff --git a/kernels/compiler_ribbon_ref.bmp b/kernels/compiler_ribbon_ref.bmp Binary files differdeleted file mode 100644 index 2225f454..00000000 --- a/kernels/compiler_ribbon_ref.bmp +++ /dev/null diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index f609cc2d..8cc8b433 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -22,7 +22,6 @@ set (utests_sources utest_error.c compiler_basic_arithmetic.cpp compiler_displacement_map_element.cpp - compiler_shader_toy.cpp compiler_mandelbrot.cpp compiler_mandelbrot_alternate.cpp compiler_box_blur_float.cpp diff --git a/utests/compiler_shader_toy.cpp b/utests/compiler_shader_toy.cpp deleted file mode 100644 index 58bcc6f7..00000000 --- a/utests/compiler_shader_toy.cpp +++ /dev/null @@ -1,87 +0,0 @@ -/* - * Copyright © 2012 Intel Corporation - * - * This library is free software; you can redistribute it and/or - * modify it under the terms of the GNU Lesser General Public - * License as published by the Free Software Foundation; either - * version 2 of the License, or (at your option) any later version. - * - * This library is distributed in the hope that it will be useful, - * but WITHOUT ANY WARRANTY; without even the implied warranty of - * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU - * Lesser General Public License for more details. - * - * You should have received a copy of the GNU Lesser General Public - * License along with this library. If not, see <http://www.gnu.org/licenses/>. - * - * Author: Benjamin Segovia <benjamin.segovia@intel.com> - */ - -/* This is a super simple wrapper for the OpenCL kernels I ported from GLSL code - * taken in Inigo's web site: - * http://www.iquilezles.org/apps/shadertoy/index.html - * - * They are pretty cool and rather complex kernels. Just the right thing to have - * something a bit more complicated and interesting than unit tests. - * - * The code here is just to wrap the common code used by all the kernels (to run - * the code and assert its correctness) - */ -#include "utest_helper.hpp" - -static const int dim = 256; - -// tricky here 'name' stands for Kernel and Reference -// 'file' stands for .cl file name and dst image name -static void run_kernel(int w, int h, const char *file, const char *name) -{ - const size_t global[2] = {size_t(w), size_t(h)}; - const size_t local[2] = {16, 1}; - const size_t sz = w * h * sizeof(char[4]); - const float fx = float(w); - const float fy = float(h); - char kernel_file[256]; - char dst_img[256]; - char ref_img[256]; - - snprintf(kernel_file, sizeof(kernel_file), "%s.cl", file); - snprintf(dst_img, sizeof(dst_img), "%s.bmp", file); - snprintf(ref_img, sizeof(ref_img), "%s_ref.bmp", name); - OCL_CALL (cl_kernel_init, kernel_file, name, SOURCE, NULL); - - OCL_CREATE_BUFFER(buf[0], 0, sz, NULL); - OCL_CALL (clSetKernelArg, kernel, 0, sizeof(cl_mem), &buf[0]); - OCL_CALL (clSetKernelArg, kernel, 1, sizeof(float), &fx); - OCL_CALL (clSetKernelArg, kernel, 2, sizeof(float), &fy); - OCL_CALL (clSetKernelArg, kernel, 3, sizeof(int), &w); - OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, 2, NULL, global, local, 0, NULL, NULL); - OCL_MAP_BUFFER(0); - int *dst = (int*) buf_data[0]; - - /* Save the image (for debug purpose) */ - cl_write_bmp(dst, w, h, dst_img); - - /* Compare with the golden image */ - OCL_CHECK_IMAGE(dst, w, h, ref_img); -} - -#define DECL_SHADER_TOY_TEST(W,H,FILE_NAME, KERNEL_NAME) \ - static void FILE_NAME(void) { run_kernel(W,H,#FILE_NAME, #KERNEL_NAME); } \ - MAKE_UTEST_FROM_FUNCTION(FILE_NAME); - -DECL_SHADER_TOY_TEST(dim,dim,compiler_clod,compiler_clod); -DECL_SHADER_TOY_TEST(dim,dim,compiler_ribbon,compiler_ribbon); -DECL_SHADER_TOY_TEST(dim,dim,compiler_nautilus,compiler_nautilus); -DECL_SHADER_TOY_TEST(dim,dim,compiler_menger_sponge_no_shadow,compiler_menger_sponge_no_shadow); -DECL_SHADER_TOY_TEST(dim,dim,compiler_julia,compiler_julia); -DECL_SHADER_TOY_TEST(dim,dim,compiler_julia_no_break,compiler_julia_no_break); -// test for function calls -DECL_SHADER_TOY_TEST(dim,dim,compiler_clod_function_call,compiler_clod); -DECL_SHADER_TOY_TEST(dim,dim,compiler_julia_function_call,compiler_julia); - -// Still issues here for LLVM 3.2 -// DECL_SHADER_TOY_TEST(dim,dim,compiler_chocolux,compiler_chocolux); -// DECL_SHADER_TOY_TEST(dim,dim,compiler_menger_sponge,compiler_menger_sponge); - -#undef DECL_SHADER_TOY_TEST - |