1.1 --- a/rt.cl	Sun Aug 08 09:51:45 2010 +0100
     1.2 +++ b/rt.cl	Mon Aug 09 04:18:21 2010 +0100
     1.3 @@ -4,6 +4,7 @@
     1.4  	int xsz, ysz;
     1.5  	int num_faces, num_lights;
     1.6  	int max_iter;
     1.7 +	float4 ambient;
     1.8  	int dbg;
     1.9  };
    1.10  
    1.11 @@ -34,6 +35,7 @@
    1.12  
    1.13  struct Ray {
    1.14  	float4 origin, dir;
    1.15 +	float energy, pad[3];
    1.16  };
    1.17  
    1.18  struct SurfPoint {
    1.19 @@ -43,14 +45,26 @@
    1.20  	global const struct Material *mat;
    1.21  };
    1.22  
    1.23 -#define EPSILON 1e-6
    1.24 +struct Scene {
    1.25 +	float4 ambient;
    1.26 +	global const struct Face *faces;
    1.27 +	int num_faces;
    1.28 +	global const struct Light *lights;
    1.29 +	int num_lights;
    1.30 +	global const struct Material *matlib;
    1.31 +};
    1.32  
    1.33 -float4 shade(struct Ray ray, struct SurfPoint sp,
    1.34 -		global const struct Light *lights, int num_lights);
    1.35 +#define MIN_ENERGY	0.001
    1.36 +#define EPSILON		1e-6
    1.37 +
    1.38 +float4 trace(struct Ray ray, struct Scene *scn);
    1.39 +float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
    1.40 +bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
    1.41  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
    1.42 +
    1.43  float4 reflect(float4 v, float4 n);
    1.44  float4 transform(float4 v, global const float *xform);
    1.45 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
    1.46 +void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
    1.47  float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
    1.48  
    1.49  kernel void render(global float4 *fb,
    1.50 @@ -64,79 +78,114 @@
    1.51  		global struct Face *outfaces)
    1.52  {
    1.53  	int idx = get_global_id(0);
    1.54 -	
    1.55 -	if(idx == 0) {
    1.56 -		for(int i=0; i<rinf->num_faces; i++) {
    1.57 -			outfaces[i] = faces[i];
    1.58 -		}
    1.59 -	}
    1.60  
    1.61 -	struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
    1.62 +	struct Scene scn;
    1.63 +	scn.ambient = rinf->ambient;
    1.64 +	scn.faces = faces;
    1.65 +	scn.num_faces = rinf->num_faces;
    1.66 +	scn.lights = lights;
    1.67 +	scn.num_lights = rinf->num_lights;
    1.68 +	scn.matlib = matlib;
    1.69  
    1.70 -	struct SurfPoint sp, sp0;
    1.71 -	sp0.t = FLT_MAX;
    1.72 -	sp0.obj = 0;
    1.73 -	
    1.74 -	int max_faces = min(rinf->num_faces, rinf->dbg);
    1.75 +	struct Ray ray = primrays[idx];
    1.76 +	transform_ray(&ray, xform, invtrans);
    1.77  
    1.78 -	for(int i=0; i<max_faces; i++) {
    1.79 -		if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
    1.80 -			sp0 = sp;
    1.81 -		}
    1.82 -	}
    1.83 -
    1.84 -	if(sp0.obj) {
    1.85 -		sp0.mat = matlib + sp0.obj->matid;
    1.86 -		fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
    1.87 -	} else {
    1.88 -		fb[idx] = (float4)(0, 0, 0, 0);
    1.89 -	}
    1.90 +	fb[idx] = trace(ray, &scn);
    1.91  }
    1.92  
    1.93 -float4 shade(struct Ray ray, struct SurfPoint sp,
    1.94 -		global const struct Light *lights, int num_lights)
    1.95 +float4 trace(struct Ray ray, struct Scene *scn)
    1.96  {
    1.97 -	float4 norm = sp.norm;
    1.98 +	float4 color;
    1.99 +	struct SurfPoint sp;
   1.100 +
   1.101 +	if(find_intersection(ray, scn, &sp)) {
   1.102 +		color = shade(ray, scn, &sp);
   1.103 +	} else {
   1.104 +		color = (float4)(0, 0, 0, 0);
   1.105 +	}
   1.106 +	return color;
   1.107 +}
   1.108 +
   1.109 +float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
   1.110 +{
   1.111 +	float4 norm = sp->norm;
   1.112  	bool entering = true;
   1.113 +	struct Material mat = *sp->mat;
   1.114  
   1.115  	if(dot(ray.dir, norm) >= 0.0) {
   1.116  		norm = -norm;
   1.117  		entering = false;
   1.118  	}
   1.119  
   1.120 -	float4 dcol = (float4)(0.07, 0.07, 0.07, 0);
   1.121 +	float4 dcol = scn->ambient * mat.kd;
   1.122  	float4 scol = (float4)(0, 0, 0, 0);
   1.123  
   1.124 -	for(int i=0; i<num_lights; i++) {
   1.125 -		float4 ldir = normalize(lights[i].pos - sp.pos);
   1.126 -		float4 vdir = -normalize(ray.dir);
   1.127 -		float4 vref = reflect(vdir, norm);
   1.128 +	for(int i=0; i<scn->num_lights; i++) {
   1.129 +		float4 ldir = scn->lights[i].pos - sp->pos;
   1.130  
   1.131 -		float diff = fmax(dot(ldir, norm), 0.0f);
   1.132 -		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
   1.133 +		struct Ray shadowray;
   1.134 +		shadowray.origin = sp->pos;
   1.135 +		shadowray.dir = ldir;
   1.136  
   1.137 -		dcol += sp.mat->kd * diff * lights[i].color;
   1.138 -		//scol += sp.mat->ks * spec * lights[i].color;
   1.139 +		if(!find_intersection(shadowray, scn, 0)) {
   1.140 +			ldir = normalize(ldir);
   1.141 +			float4 vdir = -normalize(ray.dir);
   1.142 +			float4 vref = reflect(vdir, norm);
   1.143 +
   1.144 +			float diff = fmax(dot(ldir, norm), 0.0f);
   1.145 +			dcol += mat.kd * diff * scn->lights[i].color;
   1.146 +
   1.147 +			//float spec = powr(fmax(dot(ldir, vref), 0.0f), mat.spow);
   1.148 +			//scol += mat.ks * spec * scn->lights[i].color;
   1.149 +		}
   1.150 +	}
   1.151 +
   1.152 +	float4 refl_col = mat.ks * mat.kr;
   1.153 +	float refl_coeff = (refl_col.x + refl_col.y + refl_col.z) / 3.0;
   1.154 +
   1.155 +	if(refl_coeff > MIN_ENERGY) {
   1.156 +		struct Ray refl_ray;
   1.157 +		refl_ray.origin = sp->pos;
   1.158 +		refl_ray.dir = reflect(-ray.dir, norm);
   1.159 +		refl_ray.energy *= refl_coeff;
   1.160 +
   1.161 +		scol += trace(refl_ray, scn) * refl_col;
   1.162  	}
   1.163  
   1.164  	return dcol + scol;
   1.165  }
   1.166  
   1.167 -float dot3(float4 a, float4 b)
   1.168 +
   1.169 +bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
   1.170  {
   1.171 -	return a.x * b.x + a.y * b.y + a.z * b.z;
   1.172 +	struct SurfPoint sp, sp0;
   1.173 +	sp0.t = 1.0;
   1.174 +	sp0.obj = 0;
   1.175 +
   1.176 +	for(int i=0; i<scn->num_faces; i++) {
   1.177 +		if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
   1.178 +			sp0 = sp;
   1.179 +		}
   1.180 +	}
   1.181 +
   1.182 +	if(!sp0.obj) {
   1.183 +		return false;
   1.184 +	}
   1.185 +
   1.186 +	if(spres) {
   1.187 +		*spres = sp0;
   1.188 +		spres->mat = scn->matlib + sp0.obj->matid;
   1.189 +	}
   1.190 +	return true;
   1.191  }
   1.192  
   1.193 -
   1.194 -bool intersect(struct Ray ray,
   1.195 -		global const struct Face *face,
   1.196 -		struct SurfPoint *sp)
   1.197 +bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
   1.198  {
   1.199  	float4 origin = ray.origin;
   1.200  	float4 dir = ray.dir;
   1.201  	float4 norm = face->normal;
   1.202  
   1.203 -	float ndotdir = dot3(dir, norm);
   1.204 +	float ndotdir = dot(dir, norm);
   1.205  
   1.206  	if(fabs(ndotdir) <= EPSILON) {
   1.207  		return false;
   1.208 @@ -145,7 +194,7 @@
   1.209  	float4 pt = face->v[0].pos;
   1.210  	float4 vec = pt - origin;
   1.211  
   1.212 -	float ndotvec = dot3(norm, vec);
   1.213 +	float ndotvec = dot(norm, vec);
   1.214  	float t = ndotvec / ndotdir;
   1.215  
   1.216  	if(t < EPSILON || t > 1.0) {
   1.217 @@ -153,8 +202,6 @@
   1.218  	}
   1.219  	pt = origin + dir * t;
   1.220  
   1.221 -	if(pt.w < 0.0) return false;
   1.222 -
   1.223  
   1.224  	float4 bc = calc_bary(pt, face, norm);
   1.225  	float bc_sum = bc.x + bc.y + bc.z;
   1.226 @@ -188,12 +235,10 @@
   1.227  	return res;
   1.228  }
   1.229  
   1.230 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
   1.231 +void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
   1.232  {
   1.233 -	struct Ray res;
   1.234 -	res.origin = transform(ray->origin, xform);
   1.235 -	res.dir = transform(ray->dir, invtrans);
   1.236 -	return res;
   1.237 +	ray->origin = transform(ray->origin, xform);
   1.238 +	ray->dir = transform(ray->dir, invtrans);
   1.239  }
   1.240  
   1.241  float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
   1.242 @@ -205,7 +250,7 @@
   1.243  	float4 v2 = face->v[2].pos - face->v[0].pos;
   1.244  	float4 xv1v2 = cross(v1, v2);
   1.245  
   1.246 -	float area = fabs(dot3(xv1v2, norm)) * 0.5;
   1.247 +	float area = fabs(dot(xv1v2, norm)) * 0.5;
   1.248  	if(area < EPSILON) {
   1.249  		return bc;
   1.250  	}
   1.251 @@ -219,9 +264,9 @@
   1.252  	float4 x20 = cross(pv2, pv0);
   1.253  	float4 x01 = cross(pv0, pv1);
   1.254  
   1.255 -	float a0 = fabs(dot3(x12, norm)) * 0.5;
   1.256 -	float a1 = fabs(dot3(x20, norm)) * 0.5;
   1.257 -	float a2 = fabs(dot3(x01, norm)) * 0.5;
   1.258 +	float a0 = fabs(dot(x12, norm)) * 0.5;
   1.259 +	float a1 = fabs(dot(x20, norm)) * 0.5;
   1.260 +	float a2 = fabs(dot(x01, norm)) * 0.5;
   1.261  
   1.262  	bc.x = a0 / area;
   1.263  	bc.y = a1 / area;