# HG changeset patch
# User John Tsiombikas <nuclear@member.fsf.org>
# Date 1281323901 -3600
# Node ID 9e4a28063394e654e97b4c0f03df5e7ad5f4d681
# Parent  754faf15ba36959ec882349af7eab355026c2cbc
cl compiler segfaults...

diff -r 754faf15ba36 -r 9e4a28063394 rt.cl
--- a/rt.cl	Sun Aug 08 09:51:45 2010 +0100
+++ b/rt.cl	Mon Aug 09 04:18:21 2010 +0100
@@ -4,6 +4,7 @@
 	int xsz, ysz;
 	int num_faces, num_lights;
 	int max_iter;
+	float4 ambient;
 	int dbg;
 };
 
@@ -34,6 +35,7 @@
 
 struct Ray {
 	float4 origin, dir;
+	float energy, pad[3];
 };
 
 struct SurfPoint {
@@ -43,14 +45,26 @@
 	global const struct Material *mat;
 };
 
-#define EPSILON 1e-6
+struct Scene {
+	float4 ambient;
+	global const struct Face *faces;
+	int num_faces;
+	global const struct Light *lights;
+	int num_lights;
+	global const struct Material *matlib;
+};
 
-float4 shade(struct Ray ray, struct SurfPoint sp,
-		global const struct Light *lights, int num_lights);
+#define MIN_ENERGY	0.001
+#define EPSILON		1e-6
+
+float4 trace(struct Ray ray, struct Scene *scn);
+float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
+bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
+
 float4 reflect(float4 v, float4 n);
 float4 transform(float4 v, global const float *xform);
-struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
+void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
 
 kernel void render(global float4 *fb,
@@ -64,79 +78,114 @@
 		global struct Face *outfaces)
 {
 	int idx = get_global_id(0);
-	
-	if(idx == 0) {
-		for(int i=0; i<rinf->num_faces; i++) {
-			outfaces[i] = faces[i];
-		}
-	}
 
-	struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
+	struct Scene scn;
+	scn.ambient = rinf->ambient;
+	scn.faces = faces;
+	scn.num_faces = rinf->num_faces;
+	scn.lights = lights;
+	scn.num_lights = rinf->num_lights;
+	scn.matlib = matlib;
 
-	struct SurfPoint sp, sp0;
-	sp0.t = FLT_MAX;
-	sp0.obj = 0;
-	
-	int max_faces = min(rinf->num_faces, rinf->dbg);
+	struct Ray ray = primrays[idx];
+	transform_ray(&ray, xform, invtrans);
 
-	for(int i=0; i<max_faces; i++) {
-		if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
-			sp0 = sp;
-		}
-	}
-
-	if(sp0.obj) {
-		sp0.mat = matlib + sp0.obj->matid;
-		fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
-	} else {
-		fb[idx] = (float4)(0, 0, 0, 0);
-	}
+	fb[idx] = trace(ray, &scn);
 }
 
-float4 shade(struct Ray ray, struct SurfPoint sp,
-		global const struct Light *lights, int num_lights)
+float4 trace(struct Ray ray, struct Scene *scn)
 {
-	float4 norm = sp.norm;
+	float4 color;
+	struct SurfPoint sp;
+
+	if(find_intersection(ray, scn, &sp)) {
+		color = shade(ray, scn, &sp);
+	} else {
+		color = (float4)(0, 0, 0, 0);
+	}
+	return color;
+}
+
+float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
+{
+	float4 norm = sp->norm;
 	bool entering = true;
+	struct Material mat = *sp->mat;
 
 	if(dot(ray.dir, norm) >= 0.0) {
 		norm = -norm;
 		entering = false;
 	}
 
-	float4 dcol = (float4)(0.07, 0.07, 0.07, 0);
+	float4 dcol = scn->ambient * mat.kd;
 	float4 scol = (float4)(0, 0, 0, 0);
 
-	for(int i=0; i<num_lights; i++) {
-		float4 ldir = normalize(lights[i].pos - sp.pos);
-		float4 vdir = -normalize(ray.dir);
-		float4 vref = reflect(vdir, norm);
+	for(int i=0; i<scn->num_lights; i++) {
+		float4 ldir = scn->lights[i].pos - sp->pos;
 
-		float diff = fmax(dot(ldir, norm), 0.0f);
-		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
+		struct Ray shadowray;
+		shadowray.origin = sp->pos;
+		shadowray.dir = ldir;
 
-		dcol += sp.mat->kd * diff * lights[i].color;
-		//scol += sp.mat->ks * spec * lights[i].color;
+		if(!find_intersection(shadowray, scn, 0)) {
+			ldir = normalize(ldir);
+			float4 vdir = -normalize(ray.dir);
+			float4 vref = reflect(vdir, norm);
+
+			float diff = fmax(dot(ldir, norm), 0.0f);
+			dcol += mat.kd * diff * scn->lights[i].color;
+
+			//float spec = powr(fmax(dot(ldir, vref), 0.0f), mat.spow);
+			//scol += mat.ks * spec * scn->lights[i].color;
+		}
+	}
+
+	float4 refl_col = mat.ks * mat.kr;
+	float refl_coeff = (refl_col.x + refl_col.y + refl_col.z) / 3.0;
+
+	if(refl_coeff > MIN_ENERGY) {
+		struct Ray refl_ray;
+		refl_ray.origin = sp->pos;
+		refl_ray.dir = reflect(-ray.dir, norm);
+		refl_ray.energy *= refl_coeff;
+
+		scol += trace(refl_ray, scn) * refl_col;
 	}
 
 	return dcol + scol;
 }
 
-float dot3(float4 a, float4 b)
+
+bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
 {
-	return a.x * b.x + a.y * b.y + a.z * b.z;
+	struct SurfPoint sp, sp0;
+	sp0.t = 1.0;
+	sp0.obj = 0;
+
+	for(int i=0; i<scn->num_faces; i++) {
+		if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
+			sp0 = sp;
+		}
+	}
+
+	if(!sp0.obj) {
+		return false;
+	}
+
+	if(spres) {
+		*spres = sp0;
+		spres->mat = scn->matlib + sp0.obj->matid;
+	}
+	return true;
 }
 
-
-bool intersect(struct Ray ray,
-		global const struct Face *face,
-		struct SurfPoint *sp)
+bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
 {
 	float4 origin = ray.origin;
 	float4 dir = ray.dir;
 	float4 norm = face->normal;
 
-	float ndotdir = dot3(dir, norm);
+	float ndotdir = dot(dir, norm);
 
 	if(fabs(ndotdir) <= EPSILON) {
 		return false;
@@ -145,7 +194,7 @@
 	float4 pt = face->v[0].pos;
 	float4 vec = pt - origin;
 
-	float ndotvec = dot3(norm, vec);
+	float ndotvec = dot(norm, vec);
 	float t = ndotvec / ndotdir;
 
 	if(t < EPSILON || t > 1.0) {
@@ -153,8 +202,6 @@
 	}
 	pt = origin + dir * t;
 
-	if(pt.w < 0.0) return false;
-
 
 	float4 bc = calc_bary(pt, face, norm);
 	float bc_sum = bc.x + bc.y + bc.z;
@@ -188,12 +235,10 @@
 	return res;
 }
 
-struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
+void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
 {
-	struct Ray res;
-	res.origin = transform(ray->origin, xform);
-	res.dir = transform(ray->dir, invtrans);
-	return res;
+	ray->origin = transform(ray->origin, xform);
+	ray->dir = transform(ray->dir, invtrans);
 }
 
 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
@@ -205,7 +250,7 @@
 	float4 v2 = face->v[2].pos - face->v[0].pos;
 	float4 xv1v2 = cross(v1, v2);
 
-	float area = fabs(dot3(xv1v2, norm)) * 0.5;
+	float area = fabs(dot(xv1v2, norm)) * 0.5;
 	if(area < EPSILON) {
 		return bc;
 	}
@@ -219,9 +264,9 @@
 	float4 x20 = cross(pv2, pv0);
 	float4 x01 = cross(pv0, pv1);
 
-	float a0 = fabs(dot3(x12, norm)) * 0.5;
-	float a1 = fabs(dot3(x20, norm)) * 0.5;
-	float a2 = fabs(dot3(x01, norm)) * 0.5;
+	float a0 = fabs(dot(x12, norm)) * 0.5;
+	float a1 = fabs(dot(x20, norm)) * 0.5;
+	float a2 = fabs(dot(x01, norm)) * 0.5;
 
 	bc.x = a0 / area;
 	bc.y = a1 / area;
diff -r 754faf15ba36 -r 9e4a28063394 src/mesh.cc
--- a/src/mesh.cc	Sun Aug 08 09:51:45 2010 +0100
+++ b/src/mesh.cc	Mon Aug 09 04:18:21 2010 +0100
@@ -28,6 +28,7 @@
 	CMD(KA),		\
 	CMD(KD),		\
 	CMD(KS),		\
+	CMD(KR),		\
 	CMD(NS),		\
 	CMD(NI),		\
 	CMD(D),			\
@@ -84,6 +85,7 @@
 	float shininess;	// Ns
 	float ior;			// Ni
 	float alpha;		// d, Tr
+	float refl;			// Kr (my extesnsion)
 
 	string tex_dif, tex_spec, tex_shin, tex_alpha;	// map_Kd, map_Ks, map_Ns, map_d
 	string tex_refl;	// refl -type sphere|cube file
@@ -316,7 +318,7 @@
 					mat.ks[2] = vmtl[i].specular.z;
 
 					mat.kt = 1.0 - vmtl[i].alpha;
-					mat.kr = 0.0;	// TODO
+					mat.kr = vmtl[i].refl;
 					mat.spow = vmtl[i].shininess;
 
 					matlib.push_back(mat);
@@ -487,6 +489,12 @@
 			parse_color(&mat.specular);
 			break;
 
+		case CMD_KR:
+			if((tok = strtok(0, SEP)) && is_float(tok)) {
+				mat.refl = atof(tok);
+			}
+			break;
+
 		case CMD_NS:
 			if((tok = strtok(0, SEP)) && is_float(tok)) {
 				mat.shininess = atof(tok);
diff -r 754faf15ba36 -r 9e4a28063394 src/rt.cc
--- a/src/rt.cc	Sun Aug 08 09:51:45 2010 +0100
+++ b/src/rt.cc	Mon Aug 09 04:18:21 2010 +0100
@@ -25,11 +25,13 @@
 	int xsz, ysz;
 	int num_faces, num_lights;
 	int max_iter;
+	float ambient[4];
 	int dbg;
 };
 
 struct Ray {
 	float origin[4], dir[4];
+	float energy;
 };
 
 struct Light {
@@ -45,7 +47,7 @@
 static int global_size;
 
 static Light lightlist[] = {
-	{{-10, 10, -20, 0}, {1, 1, 1, 1}}
+	{{-10, 13, -20, 0}, {1, 1, 1, 1}}
 };
 
 
@@ -55,6 +57,9 @@
 bool init_renderer(int xsz, int ysz, Scene *scn)
 {
 	// render info
+	rinf.ambient[0] = rinf.ambient[1] = rinf.ambient[2] = 0.075;
+	rinf.ambient[3] = 0.0;
+
 	rinf.xsz = xsz;
 	rinf.ysz = ysz;
 	rinf.num_faces = scn->get_num_faces();
@@ -98,6 +103,10 @@
 		return false;
 	}
 
+	if(!prog->build()) {
+		return false;
+	}
+
 	delete [] prim_rays;
 
 	global_size = xsz * ysz;
@@ -250,7 +259,7 @@
 	py *= 100.0;
 	pz *= 100.0;
 
-	Ray ray = {{0, 0, 0, 1}, {px, py, -pz, 1}};
+	Ray ray = {{0, 0, 0, 1}, {px, py, -pz, 1}, 1.0};
 	return ray;
 }
 
diff -r 754faf15ba36 -r 9e4a28063394 test.cl
--- a/test.cl	Sun Aug 08 09:51:45 2010 +0100
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,6 +0,0 @@
-__kernel void test(__global int *dst, __global const int *src, __global const int foo)
-{
-	int idx = get_global_id(0);
-
-	dst[idx] = src[idx] * foo;
-}