[Beignet] [PATCH 3/3] utests: add test cases for function call.
Ruiling Song
ruiling.song at intel.com
Wed Oct 16 00:38:08 PDT 2013
Signed-off-by: Ruiling Song <ruiling.song at intel.com>
---
kernels/compiler_clod_function_call.cl | 91 ++++++++++++++++++++
kernels/compiler_julia_function_call.cl | 142 +++++++++++++++++++++++++++++++
utests/compiler_shader_toy.cpp | 33 ++++---
3 files changed, 252 insertions(+), 14 deletions(-)
create mode 100644 kernels/compiler_clod_function_call.cl
create mode 100644 kernels/compiler_julia_function_call.cl
diff --git a/kernels/compiler_clod_function_call.cl b/kernels/compiler_clod_function_call.cl
new file mode 100644
index 0000000..ecfac46
--- /dev/null
+++ b/kernels/compiler_clod_function_call.cl
@@ -0,0 +1,91 @@
+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
+
+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)
+
+#define time 1.f
+
+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
+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_julia_function_call.cl b/kernels/compiler_julia_function_call.cl
new file mode 100644
index 0000000..7b3aa46
--- /dev/null
+++ b/kernels/compiler_julia_function_call.cl
@@ -0,0 +1,142 @@
+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/utests/compiler_shader_toy.cpp b/utests/compiler_shader_toy.cpp
index ead9120..58bcc6f 100644
--- a/utests/compiler_shader_toy.cpp
+++ b/utests/compiler_shader_toy.cpp
@@ -31,7 +31,9 @@
static const int dim = 256;
-static void run_kernel(int w, int h, const char *name)
+// 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};
@@ -42,8 +44,8 @@ static void run_kernel(int w, int h, const char *name)
char dst_img[256];
char ref_img[256];
- snprintf(kernel_file, sizeof(kernel_file), "%s.cl", name);
- snprintf(dst_img, sizeof(dst_img), "%s.bmp", name);
+ 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);
@@ -63,20 +65,23 @@ static void run_kernel(int w, int h, const char *name)
OCL_CHECK_IMAGE(dst, w, h, ref_img);
}
-#define DECL_SHADER_TOY_TEST(W,H,NAME) \
- static void NAME(void) { run_kernel(W,H,#NAME); } \
- MAKE_UTEST_FROM_FUNCTION(NAME);
+#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);
-DECL_SHADER_TOY_TEST(dim,dim,compiler_ribbon);
-DECL_SHADER_TOY_TEST(dim,dim,compiler_nautilus);
-DECL_SHADER_TOY_TEST(dim,dim,compiler_menger_sponge_no_shadow);
-DECL_SHADER_TOY_TEST(dim,dim,compiler_julia);
-DECL_SHADER_TOY_TEST(dim,dim,compiler_julia_no_break);
+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);
-// DECL_SHADER_TOY_TEST(dim,dim,compiler_menger_sponge);
+// 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
--
1.7.9.5
More information about the Beignet
mailing list