nuclear@2: struct RendInfo {
nuclear@2: 	int xsz, ysz;
nuclear@3: 	int num_sph, num_lights;
nuclear@2: 	int max_iter;
nuclear@2: };
nuclear@2: 
nuclear@2: struct Sphere {
nuclear@2: 	float4 pos;
nuclear@5: 	float4 kd, ks;
nuclear@2: 	float radius;
nuclear@4: 	float spow, kr, kt;
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@2: 	float3 pos, norm;
nuclear@4: 	global const struct Sphere *obj;
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@4: bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
nuclear@5: float3 reflect(float3 v, float3 n);
nuclear@2: 
nuclear@4: 
nuclear@4: kernel void render(global float4 *fb,
nuclear@4: 		global const struct RendInfo *rinf,
nuclear@4: 		global const struct Sphere *sphlist,
nuclear@4: 		global const struct Light *lights,
nuclear@4: 		global const struct Ray *primrays)
nuclear@2: {
nuclear@2: 	int idx = get_global_id(0);
nuclear@2: 
nuclear@2: 	struct Ray ray = primrays[idx];
nuclear@4: 	struct SurfPoint sp, sp0;
nuclear@2: 
nuclear@4: 	sp0.t = FLT_MAX;
nuclear@5: 	sp0.obj = 0;
nuclear@4: 
nuclear@4: 	for(int i=0; i<rinf->num_sph; i++) {
nuclear@5: 		if(intersect(ray, sphlist + i, &sp) && sp.t < sp0.t) {
nuclear@4: 			sp0 = sp;
nuclear@4: 		}
nuclear@2: 	}
nuclear@3: 
nuclear@5: 	if(sp0.obj) {
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@5: 	float3 dcol = (float3)(0, 0, 0);
nuclear@5: 	float3 scol = (float3)(0, 0, 0);
nuclear@5: 
nuclear@5: 	for(int i=0; i<num_lights; i++) {
nuclear@5: 		float3 ldir = normalize(lights[i].pos.xyz - sp.pos);
nuclear@5: 		float3 vdir = -normalize(ray.dir.xyz);
nuclear@5: 		float3 vref = reflect(vdir, sp.norm);
nuclear@5: 
nuclear@5: 		float diff = fmax(dot(ldir, sp.norm), 0.0f);
nuclear@5: 		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.obj->spow);
nuclear@5: 
nuclear@5: 		dcol += sp.obj->kd.xyz * diff * lights[i].color.xyz;
nuclear@5: 		scol += sp.obj->ks.xyz * spec * lights[i].color.xyz;
nuclear@5: 	}
nuclear@5: 
nuclear@5: 	return (float4)(dcol + scol, 1.0f);
nuclear@2: }
nuclear@2: 
nuclear@2: bool intersect(struct Ray ray,
nuclear@4: 		global const struct Sphere *sph,
nuclear@2: 		struct SurfPoint *sp)
nuclear@2: {
nuclear@2: 	float3 dir = ray.dir.xyz;
nuclear@2: 	float3 orig = ray.origin.xyz;
nuclear@2: 	float3 spos = sph->pos.xyz;
nuclear@2: 
nuclear@2: 	float a = dot(dir, dir);
nuclear@2: 	float b = 2.0 * dir.x * (orig.x - spos.x) +
nuclear@2: 		2.0 * dir.y * (orig.y - spos.y) +
nuclear@2: 		2.0 * dir.z * (orig.z - spos.z);
nuclear@2: 	float c = dot(spos, spos) + dot(orig, orig) +
nuclear@2: 		2.0 * dot(-spos, orig) - sph->radius * sph->radius;
nuclear@2: 
nuclear@2: 	float d = b * b - 4.0 * a * c;
nuclear@2: 	if(d < 0.0) return false;
nuclear@2: 
nuclear@2: 	float sqrt_d = sqrt(d);
nuclear@2: 	float t1 = (-b + sqrt_d) / (2.0 * a);
nuclear@2: 	float t2 = (-b - sqrt_d) / (2.0 * a);
nuclear@2: 
nuclear@2: 	if(t1 < EPSILON) t1 = t2;
nuclear@2: 	if(t2 < EPSILON) t2 = t1;
nuclear@2: 	float t = t1 < t2 ? t1 : t2;
nuclear@2: 
nuclear@2: 	if(t < EPSILON || t > 1.0) {
nuclear@2: 		return false;
nuclear@2: 	}
nuclear@2: 
nuclear@2: 	sp->t = t;
nuclear@2: 	sp->pos = orig + dir * sp->t;
nuclear@2: 	sp->norm = (sp->pos - spos) / sph->radius;
nuclear@5: 	sp->obj = sph;
nuclear@2: 	return true;
nuclear@2: }
nuclear@5: 
nuclear@5: float3 reflect(float3 v, float3 n)
nuclear@5: {
nuclear@5: 	return 2.0f * dot(v, n) * n - v;
nuclear@5: }