1.1 --- a/rt.cl	Sat Jul 31 22:23:57 2010 +0100
     1.2 +++ b/rt.cl	Tue Aug 03 13:06:59 2010 +0100
     1.3 @@ -1,3 +1,5 @@
     1.4 +/* vim: set ft=opencl:ts=4:sw=4 */
     1.5 +
     1.6  struct RendInfo {
     1.7  	int xsz, ysz;
     1.8  	int num_faces, num_lights;
     1.9 @@ -7,19 +9,22 @@
    1.10  struct Vertex {
    1.11  	float4 pos;
    1.12  	float4 normal;
    1.13 -	float2 tex;
    1.14 +	float4 tex;
    1.15 +	float4 padding;
    1.16  };
    1.17  
    1.18  struct Face {
    1.19  	struct Vertex v[3];
    1.20  	float4 normal;
    1.21  	int matid;
    1.22 +	int padding[3];
    1.23  };
    1.24  
    1.25  struct Material {
    1.26  	float4 kd, ks;
    1.27  	float kr, kt;
    1.28  	float spow;
    1.29 +	float padding;
    1.30  };
    1.31  
    1.32  struct Light {
    1.33 @@ -32,7 +37,7 @@
    1.34  
    1.35  struct SurfPoint {
    1.36  	float t;
    1.37 -	float4 pos, norm;
    1.38 +	float4 pos, norm, dbg;
    1.39  	global const struct Face *obj;
    1.40  	global const struct Material *mat;
    1.41  };
    1.42 @@ -44,9 +49,8 @@
    1.43  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
    1.44  float4 reflect(float4 v, float4 n);
    1.45  float4 transform(float4 v, global const float *xform);
    1.46 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform);
    1.47 -float4 calc_bary(float4 pt, global const struct Face *face);
    1.48 -
    1.49 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
    1.50 +float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
    1.51  
    1.52  kernel void render(global float4 *fb,
    1.53  		global const struct RendInfo *rinf,
    1.54 @@ -54,11 +58,12 @@
    1.55  		global const struct Material *matlib,
    1.56  		global const struct Light *lights,
    1.57  		global const struct Ray *primrays,
    1.58 -		global const float *xform)
    1.59 +		global const float *xform,
    1.60 +		global const float *invtrans)
    1.61  {
    1.62  	int idx = get_global_id(0);
    1.63  
    1.64 -	struct Ray ray = transform_ray(primrays + idx, xform);
    1.65 +	struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
    1.66  
    1.67  	struct SurfPoint sp, sp0;
    1.68  	sp0.t = FLT_MAX;
    1.69 @@ -81,61 +86,86 @@
    1.70  float4 shade(struct Ray ray, struct SurfPoint sp,
    1.71  		global const struct Light *lights, int num_lights)
    1.72  {
    1.73 +	float4 norm = sp.norm;
    1.74 +	bool entering = true;
    1.75 +
    1.76 +	if(dot(ray.dir, norm) >= 0.0) {
    1.77 +		norm = -norm;
    1.78 +		entering = false;
    1.79 +	}
    1.80 +
    1.81  	float4 dcol = (float4)(0, 0, 0, 0);
    1.82  	float4 scol = (float4)(0, 0, 0, 0);
    1.83  
    1.84  	for(int i=0; i<num_lights; i++) {
    1.85  		float4 ldir = normalize(lights[i].pos - sp.pos);
    1.86  		float4 vdir = -normalize(ray.dir);
    1.87 -		float4 vref = reflect(vdir, sp.norm);
    1.88 +		float4 vref = reflect(vdir, norm);
    1.89  
    1.90 -		float diff = fmax(dot(ldir, sp.norm), 0.0f);
    1.91 +		float diff = fmax(dot(ldir, norm), 0.0f);
    1.92  		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
    1.93  
    1.94  		dcol += sp.mat->kd * diff * lights[i].color;
    1.95 -		scol += sp.mat->ks * spec * lights[i].color;
    1.96 +		//scol += sp.mat->ks * spec * lights[i].color;
    1.97  	}
    1.98  
    1.99  	return dcol + scol;
   1.100  }
   1.101  
   1.102 +float dot3(float4 a, float4 b)
   1.103 +{
   1.104 +	return a.x * b.x + a.y * b.y + a.z * b.z;
   1.105 +}
   1.106 +
   1.107 +
   1.108  bool intersect(struct Ray ray,
   1.109  		global const struct Face *face,
   1.110  		struct SurfPoint *sp)
   1.111  {
   1.112 -	float ndotdir = dot(face->normal, ray.dir);
   1.113 +	float4 origin = ray.origin;
   1.114 +	float4 dir = ray.dir;
   1.115 +	float4 norm = face->normal;
   1.116 +
   1.117 +	float ndotdir = dot3(dir, norm);
   1.118 +
   1.119  	if(fabs(ndotdir) <= EPSILON) {
   1.120  		return false;
   1.121  	}
   1.122  
   1.123  	float4 pt = face->v[0].pos;
   1.124 -	float4 vec = pt - ray.origin;
   1.125 +	float4 vec = pt - origin;
   1.126  
   1.127 -	float ndotvec = dot(face->normal, vec);
   1.128 +	float ndotvec = dot3(norm, vec);
   1.129  	float t = ndotvec / ndotdir;
   1.130  
   1.131  	if(t < EPSILON || t > 1.0) {
   1.132  		return false;
   1.133  	}
   1.134 -	pt = ray.origin + ray.dir * t;
   1.135 +	pt = origin + dir * t;
   1.136  
   1.137 -	float4 bc = calc_bary(pt, face);
   1.138 +	if(pt.w < 0.0) return false;
   1.139 +
   1.140 +
   1.141 +	float4 bc = calc_bary(pt, face, norm);
   1.142  	float bc_sum = bc.x + bc.y + bc.z;
   1.143  
   1.144 -	if(bc_sum < -EPSILON || bc_sum > 1.0) {
   1.145 +	if(bc_sum < 0.0 || bc_sum > 1.0 + EPSILON) {
   1.146  		return false;
   1.147 +		bc *= 1.2;
   1.148  	}
   1.149  
   1.150  	sp->t = t;
   1.151  	sp->pos = pt;
   1.152 -	sp->norm = face->normal;
   1.153 +	sp->norm = norm;
   1.154  	sp->obj = face;
   1.155 +	sp->dbg = bc;
   1.156  	return true;
   1.157  }
   1.158  
   1.159  float4 reflect(float4 v, float4 n)
   1.160  {
   1.161 -	return 2.0f * dot(v, n) * n - v;
   1.162 +	float4 res = 2.0f * dot(v, n) * n - v;
   1.163 +	return res;
   1.164  }
   1.165  
   1.166  float4 transform(float4 v, global const float *xform)
   1.167 @@ -144,33 +174,28 @@
   1.168  	res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
   1.169  	res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
   1.170  	res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
   1.171 -	res.w = 1.0;
   1.172 +	res.w = 0.0;
   1.173  	return res;
   1.174  }
   1.175  
   1.176 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform)
   1.177 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
   1.178  {
   1.179  	struct Ray res;
   1.180 -	float rot[16];
   1.181 -
   1.182 -	for(int i=0; i<16; i++) {
   1.183 -		rot[i] = xform[i];
   1.184 -	}
   1.185 -	rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f;
   1.186 -	rot[15] = 1.0f;
   1.187 -
   1.188  	res.origin = transform(ray->origin, xform);
   1.189 -	res.dir = transform(ray->dir, xform);
   1.190 +	res.dir = transform(ray->dir, invtrans);
   1.191  	return res;
   1.192  }
   1.193  
   1.194 -float4 calc_bary(float4 pt, global const struct Face *face)
   1.195 +float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
   1.196  {
   1.197 -	float4 bc = {0, 0, 0, 0};
   1.198 +	float4 bc = (float4)(0, 0, 0, 0);
   1.199  
   1.200 -	float4 vi = face->v[1].pos - face->v[0].pos;
   1.201 -	float4 vj = face->v[2].pos - face->v[0].pos;
   1.202 -	float area = fabs(dot(cross(vi, vj), face->normal) / 2.0);
   1.203 +	// calculate area of the whole triangle
   1.204 +	float4 v1 = face->v[1].pos - face->v[0].pos;
   1.205 +	float4 v2 = face->v[2].pos - face->v[0].pos;
   1.206 +	float4 xv1v2 = cross(v1, v2);
   1.207 +
   1.208 +	float area = fabs(dot3(xv1v2, norm)) * 0.5;
   1.209  	if(area < EPSILON) {
   1.210  		return bc;
   1.211  	}
   1.212 @@ -179,10 +204,14 @@
   1.213  	float4 pv1 = face->v[1].pos - pt;
   1.214  	float4 pv2 = face->v[2].pos - pt;
   1.215  
   1.216 -	// calculate the areas of each sub-triangle
   1.217 -	float a0 = fabs(dot(cross(pv1, pv2), face->normal) / 2.0);
   1.218 -	float a1 = fabs(dot(cross(pv2, pv0), face->normal) / 2.0);
   1.219 -	float a2 = fabs(dot(cross(pv0, pv1), face->normal) / 2.0);
   1.220 +	// calculate the area of each sub-triangle
   1.221 +	float4 x12 = cross(pv1, pv2);
   1.222 +	float4 x20 = cross(pv2, pv0);
   1.223 +	float4 x01 = cross(pv0, pv1);
   1.224 +
   1.225 +	float a0 = fabs(dot3(x12, norm)) * 0.5;
   1.226 +	float a1 = fabs(dot3(x20, norm)) * 0.5;
   1.227 +	float a2 = fabs(dot3(x01, norm)) * 0.5;
   1.228  
   1.229  	bc.x = a0 / area;
   1.230  	bc.y = a1 / area;