nuclear@2: struct RendInfo {
nuclear@2: 	int xsz, ysz;
nuclear@9: 	int num_faces, num_lights;
nuclear@2: 	int max_iter;
nuclear@2: };
nuclear@2: 
nuclear@9: struct Vertex {
nuclear@2: 	float4 pos;
nuclear@9: 	float4 normal;
nuclear@9: 	float2 tex;
nuclear@9: };
nuclear@9: 
nuclear@9: struct Face {
nuclear@9: 	struct Vertex v[3];
nuclear@9: 	float4 normal;
nuclear@9: 	int matid;
nuclear@9: };
nuclear@9: 
nuclear@9: struct Material {
nuclear@5: 	float4 kd, ks;
nuclear@9: 	float kr, kt;
nuclear@9: 	float spow;
nuclear@2: };
nuclear@2: 
nuclear@3: struct Light {
nuclear@3: 	float4 pos, color;
nuclear@3: };
nuclear@3: 
nuclear@2: struct Ray {
nuclear@2: 	float4 origin, dir;
nuclear@2: };
nuclear@2: 
nuclear@2: struct SurfPoint {
nuclear@2: 	float t;
nuclear@8: 	float4 pos, norm;
nuclear@9: 	global const struct Face *obj;
nuclear@9: 	global const struct Material *mat;
nuclear@2: };
nuclear@2: 
nuclear@2: #define EPSILON 1e-6
nuclear@2: 
nuclear@5: float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5: 		global const struct Light *lights, int num_lights);
nuclear@9: bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@8: float4 reflect(float4 v, float4 n);
nuclear@8: float4 transform(float4 v, global const float *xform);
nuclear@8: struct Ray transform_ray(global const struct Ray *ray, global const float *xform);
nuclear@9: float4 calc_bary(float4 pt, global const struct Face *face);
nuclear@2: 
nuclear@4: 
nuclear@4: kernel void render(global float4 *fb,
nuclear@4: 		global const struct RendInfo *rinf,
nuclear@9: 		global const struct Face *faces,
nuclear@9: 		global const struct Material *matlib,
nuclear@4: 		global const struct Light *lights,
nuclear@7: 		global const struct Ray *primrays,
nuclear@8: 		global const float *xform)
nuclear@2: {
nuclear@2: 	int idx = get_global_id(0);
nuclear@2: 
nuclear@8: 	struct Ray ray = transform_ray(primrays + idx, xform);
nuclear@8: 
nuclear@4: 	struct SurfPoint sp, sp0;
nuclear@4: 	sp0.t = FLT_MAX;
nuclear@5: 	sp0.obj = 0;
nuclear@4: 
nuclear@9: 	for(int i=0; i<rinf->num_faces; i++) {
nuclear@9: 		if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
nuclear@4: 			sp0 = sp;
nuclear@4: 		}
nuclear@2: 	}
nuclear@3: 
nuclear@5: 	if(sp0.obj) {
nuclear@9: 		sp0.mat = matlib + sp0.obj->matid;
nuclear@5: 		fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
nuclear@5: 	} else {
nuclear@5: 		fb[idx] = (float4)(0, 0, 0, 0);
nuclear@5: 	}
nuclear@4: }
nuclear@4: 
nuclear@5: float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5: 		global const struct Light *lights, int num_lights)
nuclear@4: {
nuclear@8: 	float4 dcol = (float4)(0, 0, 0, 0);
nuclear@8: 	float4 scol = (float4)(0, 0, 0, 0);
nuclear@5: 
nuclear@5: 	for(int i=0; i<num_lights; i++) {
nuclear@8: 		float4 ldir = normalize(lights[i].pos - sp.pos);
nuclear@8: 		float4 vdir = -normalize(ray.dir);
nuclear@8: 		float4 vref = reflect(vdir, sp.norm);
nuclear@5: 
nuclear@5: 		float diff = fmax(dot(ldir, sp.norm), 0.0f);
nuclear@9: 		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
nuclear@5: 
nuclear@9: 		dcol += sp.mat->kd * diff * lights[i].color;
nuclear@9: 		scol += sp.mat->ks * spec * lights[i].color;
nuclear@5: 	}
nuclear@5: 
nuclear@8: 	return dcol + scol;
nuclear@2: }
nuclear@2: 
nuclear@2: bool intersect(struct Ray ray,
nuclear@9: 		global const struct Face *face,
nuclear@2: 		struct SurfPoint *sp)
nuclear@2: {
nuclear@9: 	float ndotdir = dot(face->normal, ray.dir);
nuclear@9: 	if(fabs(ndotdir) <= EPSILON) {
nuclear@9: 		return false;
nuclear@9: 	}
nuclear@2: 
nuclear@9: 	float4 pt = face->v[0].pos;
nuclear@9: 	float4 vec = pt - ray.origin;
nuclear@2: 
nuclear@9: 	float ndotvec = dot(face->normal, vec);
nuclear@9: 	float t = ndotvec / ndotdir;
nuclear@2: 
nuclear@2: 	if(t < EPSILON || t > 1.0) {
nuclear@2: 		return false;
nuclear@2: 	}
nuclear@9: 	pt = ray.origin + ray.dir * t;
nuclear@9: 
nuclear@9: 	float4 bc = calc_bary(pt, face);
nuclear@9: 	float bc_sum = bc.x + bc.y + bc.z;
nuclear@9: 
nuclear@9: 	if(bc_sum < -EPSILON || bc_sum > 1.0) {
nuclear@9: 		return false;
nuclear@9: 	}
nuclear@2: 
nuclear@2: 	sp->t = t;
nuclear@9: 	sp->pos = pt;
nuclear@9: 	sp->norm = face->normal;
nuclear@9: 	sp->obj = face;
nuclear@2: 	return true;
nuclear@2: }
nuclear@5: 
nuclear@8: float4 reflect(float4 v, float4 n)
nuclear@5: {
nuclear@5: 	return 2.0f * dot(v, n) * n - v;
nuclear@5: }
nuclear@8: 
nuclear@8: float4 transform(float4 v, global const float *xform)
nuclear@8: {
nuclear@8: 	float4 res;
nuclear@8: 	res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8: 	res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8: 	res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@8: 	res.w = 1.0;
nuclear@8: 	return res;
nuclear@8: }
nuclear@8: 
nuclear@8: struct Ray transform_ray(global const struct Ray *ray, global const float *xform)
nuclear@8: {
nuclear@8: 	struct Ray res;
nuclear@8: 	float rot[16];
nuclear@8: 
nuclear@8: 	for(int i=0; i<16; i++) {
nuclear@8: 		rot[i] = xform[i];
nuclear@8: 	}
nuclear@8: 	rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f;
nuclear@8: 	rot[15] = 1.0f;
nuclear@8: 
nuclear@8: 	res.origin = transform(ray->origin, xform);
nuclear@8: 	res.dir = transform(ray->dir, xform);
nuclear@8: 	return res;
nuclear@8: }
nuclear@9: 
nuclear@9: float4 calc_bary(float4 pt, global const struct Face *face)
nuclear@9: {
nuclear@9: 	float4 bc = {0, 0, 0, 0};
nuclear@9: 
nuclear@9: 	float4 vi = face->v[1].pos - face->v[0].pos;
nuclear@9: 	float4 vj = face->v[2].pos - face->v[0].pos;
nuclear@9: 	float area = fabs(dot(cross(vi, vj), face->normal) / 2.0);
nuclear@9: 	if(area < EPSILON) {
nuclear@9: 		return bc;
nuclear@9: 	}
nuclear@9: 
nuclear@9: 	float4 pv0 = face->v[0].pos - pt;
nuclear@9: 	float4 pv1 = face->v[1].pos - pt;
nuclear@9: 	float4 pv2 = face->v[2].pos - pt;
nuclear@9: 
nuclear@9: 	// calculate the areas of each sub-triangle
nuclear@9: 	float a0 = fabs(dot(cross(pv1, pv2), face->normal) / 2.0);
nuclear@9: 	float a1 = fabs(dot(cross(pv2, pv0), face->normal) / 2.0);
nuclear@9: 	float a2 = fabs(dot(cross(pv0, pv1), face->normal) / 2.0);
nuclear@9: 
nuclear@9: 	bc.x = a0 / area;
nuclear@9: 	bc.y = a1 / area;
nuclear@9: 	bc.z = a2 / area;
nuclear@9: 	return bc;
nuclear@9: }