1.1 --- a/rt.cl	Tue Aug 17 20:35:00 2010 +0100
     1.2 +++ b/rt.cl	Sat Aug 21 03:42:49 2010 +0100
     1.3 @@ -5,6 +5,7 @@
     1.4  	int xsz, ysz;
     1.5  	int num_faces, num_lights;
     1.6  	int max_iter;
     1.7 +	int kd_depth;
     1.8  };
     1.9  
    1.10  struct Vertex {
    1.11 @@ -50,6 +51,18 @@
    1.12  	global const struct Light *lights;
    1.13  	int num_lights;
    1.14  	global const struct Material *matlib;
    1.15 +	global const struct KDNode *kdtree;
    1.16 +};
    1.17 +
    1.18 +struct AABBox {
    1.19 +	float4 min, max;
    1.20 +};
    1.21 +
    1.22 +struct KDNode {
    1.23 +	AABBox aabb;
    1.24 +	int face_idx[32];
    1.25 +	int num_faces;
    1.26 +	int padding[3];
    1.27  };
    1.28  
    1.29  #define MIN_ENERGY	0.001
    1.30 @@ -58,6 +71,7 @@
    1.31  float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
    1.32  bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
    1.33  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
    1.34 +bool intersect_aabb(struct Ray ray, struct AABBox aabb);
    1.35  
    1.36  float4 reflect(float4 v, float4 n);
    1.37  float4 transform(float4 v, global const float *xform);
    1.38 @@ -72,7 +86,8 @@
    1.39  		global const struct Light *lights,
    1.40  		global const struct Ray *primrays,
    1.41  		global const float *xform,
    1.42 -		global const float *invtrans)
    1.43 +		global const float *invtrans,
    1.44 +		global const struct KDNode *kdtree)
    1.45  {
    1.46  	int idx = get_global_id(0);
    1.47  
    1.48 @@ -146,8 +161,12 @@
    1.49  	return dcol + scol;
    1.50  }
    1.51  
    1.52 +bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
    1.53 +{
    1.54 +	return false;
    1.55 +}
    1.56  
    1.57 -bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
    1.58 +/*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
    1.59  {
    1.60  	struct SurfPoint sp, sp0;
    1.61  	sp0.t = 1.0;
    1.62 @@ -168,7 +187,7 @@
    1.63  		spres->mat = scn->matlib[sp0.obj->matid];
    1.64  	}
    1.65  	return true;
    1.66 -}
    1.67 +}*/
    1.68  
    1.69  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
    1.70  {
    1.71 @@ -210,6 +229,44 @@
    1.72  	return true;
    1.73  }
    1.74  
    1.75 +bool intersect_aabb(struct Ray ray, struct AABBox aabb)
    1.76 +{
    1.77 +	if(ray.origin.x >= aabb.min.x && ray.origin.y >= aabb.min.y && ray.origin.z >= aabb.min.z &&
    1.78 +			ray.origin.x < aabb.max.x && ray.origin.y < aabb.max.y && ray.origin.z < aabb.max.z) {
    1.79 +		return true;
    1.80 +	}
    1.81 +
    1.82 +	float4 bbox[2] = {aabb.min, aabb.max};
    1.83 +
    1.84 +	int xsign = (int)(ray.dir.x < 0.0);
    1.85 +	float invdirx = 1.0 / ray.dir.x;
    1.86 +	float tmin = (bbox[xsign].x - ray.origin.x) * invdirx;
    1.87 +	float tmax = (bbox[1 - xsign].x - ray.origin.x) * invdirx;
    1.88 +
    1.89 +	int ysign = (int)(ray.dir.y < 0.0);
    1.90 +	float invdiry = 1.0 / ray.dir.y;
    1.91 +	float tymin = (bbox[ysign].y - ray.origin.y) * invdiry;
    1.92 +	float tymax = (bbox[1 - ysign].y - ray.origin.y) * invdiry;
    1.93 +
    1.94 +	if(tmin > tymax || tymin > tmax) {
    1.95 +		return false;
    1.96 +	}
    1.97 +
    1.98 +	if(tymin > tmin) tmin = tymin;
    1.99 +	if(tymax < tmax) tmax = tymax;
   1.100 +
   1.101 +	int zsign = (int)(ray.dir.z < 0.0);
   1.102 +	float invdirz = 1.0 / ray.dir.z;
   1.103 +	float tzmin = (bbox[zsign].z - ray.origin.z) * invdirz;
   1.104 +	float tzmax = (bbox[1 - zsign].z - ray.origin.z) * invdirz;
   1.105 +
   1.106 +	if(tmin > tzmax || tzmin > tmax) {
   1.107 +		return false;
   1.108 +	}
   1.109 +
   1.110 +	return tmin < t1 && tmax > t0;
   1.111 +}
   1.112 +
   1.113  float4 reflect(float4 v, float4 n)
   1.114  {
   1.115  	return 2.0f * dot(v, n) * n - v;