1.1 --- a/rt.cl	Tue Aug 24 05:47:04 2010 +0100
     1.2 +++ b/rt.cl	Thu Aug 26 20:24:07 2010 +0100
     1.3 @@ -62,7 +62,8 @@
     1.4  	struct AABBox aabb;
     1.5  	int face_idx[32];
     1.6  	int num_faces;
     1.7 -	int padding[3];
     1.8 +	int left, right;
     1.9 +	int padding;
    1.10  };
    1.11  
    1.12  #define MIN_ENERGY	0.001
    1.13 @@ -117,9 +118,9 @@
    1.14  			ray.origin = sp.pos;
    1.15  			ray.dir = reflect(-ray.dir, sp.norm);
    1.16  
    1.17 -			energy *= sp.mat.ks * sp.mat.kr;
    1.18 +			energy *= refl_col;
    1.19  		} else {
    1.20 -			iter = INT_MAX - 1;	// to break out of the loop
    1.21 +			break;
    1.22  		}
    1.23  	}
    1.24  
    1.25 @@ -171,7 +172,7 @@
    1.26  
    1.27  	int idxstack[STACK_SIZE];
    1.28  	int top = 0;			// points after the topmost element of the stack
    1.29 -	idxstack[top++] = 1;	// root at tree[1] (heap)
    1.30 +	idxstack[top++] = 0;	// root at tree[0]
    1.31  
    1.32  	while(top > 0) {
    1.33  		int idx = idxstack[--top];	// remove this index from the stack and process it
    1.34 @@ -186,7 +187,7 @@
    1.35  		}*/
    1.36  
    1.37  		if(intersect_aabb(ray, node->aabb)) {
    1.38 -			if(node->num_faces >= 0) {
    1.39 +			if(node->left == -1) {
    1.40  				// leaf node... check each face in turn and update the nearest intersection as needed
    1.41  				for(int i=0; i<node->num_faces; i++) {
    1.42  					struct SurfPoint spt;
    1.43 @@ -199,10 +200,10 @@
    1.44  			} else {
    1.45  				// internal node... recurse to the children
    1.46  				/*if(get_global_id(0) == 0) {
    1.47 -					printf("pushing %d's children %d and %d\n", idx, idx * 2, idx * 2 + 1);
    1.48 +					printf("pushing %d's children %d and %d\n", idx, node->left, node->right);
    1.49  				}*/
    1.50 -				idxstack[top++] = idx * 2;
    1.51 -				idxstack[top++] = idx * 2 + 1;
    1.52 +				idxstack[top++] = node->left;
    1.53 +				idxstack[top++] = node->right;
    1.54  			}
    1.55  		}
    1.56  	}

     2.1 --- a/src/rt.cc	Tue Aug 24 05:47:04 2010 +0100
     2.2 +++ b/src/rt.cc	Thu Aug 26 20:24:07 2010 +0100
     2.3 @@ -102,7 +102,7 @@
     2.4  	prog->set_arg_buffer(KARG_PRIM_RAYS, ARG_RD, xsz * ysz * sizeof *prim_rays, prim_rays);
     2.5  	prog->set_arg_buffer(KARG_XFORM, ARG_RD, 16 * sizeof(float));
     2.6  	prog->set_arg_buffer(KARG_INVTRANS_XFORM, ARG_RD, 16 * sizeof(float));
     2.7 -	prog->set_arg_buffer(KARG_KDTREE, ARG_RD, scn->get_kdtree_buffer_size(), kdbuf);
     2.8 +	prog->set_arg_buffer(KARG_KDTREE, ARG_RD, scn->get_num_kdnodes() * sizeof *kdbuf, kdbuf);
     2.9  
    2.10  	if(prog->get_num_args() < NUM_KERNEL_ARGS) {
    2.11  		return false;

     3.1 --- a/src/scene.cc	Tue Aug 24 05:47:04 2010 +0100
     3.2 +++ b/src/scene.cc	Thu Aug 26 20:24:07 2010 +0100
     3.3 @@ -1,15 +1,18 @@
     3.4 +#include <stdlib.h>
     3.5  #include <math.h>
     3.6  #include <float.h>
     3.7  #include <assert.h>
     3.8 +#include <map>
     3.9  #include "scene.h"
    3.10  #include "ogl.h"
    3.11  
    3.12  
    3.13 +static void flatten_kdtree(const KDNode *node, KDNodeGPU *kdbuf, int *count, std::map<const KDNode*, int> *flatmap);
    3.14 +static void fix_child_links(const KDNode *node, KDNodeGPU *kdbuf, const std::map<const KDNode*, int> &flatmap);
    3.15  static void draw_kdtree(const KDNode *node, int level = 0);
    3.16  static bool build_kdtree(KDNode *kd, const Face *faces, int level = 0);
    3.17  static float eval_cost(const Face *faces, const int *face_idx, int num_faces, const AABBox &aabb, int axis);
    3.18  static void free_kdtree(KDNode *node);
    3.19 -static void kdtree_gpu_flatten(KDNodeGPU *kdbuf, int idx, const KDNode *node);
    3.20  static void print_item_counts(const KDNode *node, int level);
    3.21  
    3.22  
    3.23 @@ -123,6 +126,11 @@
    3.24  	return (int)matlib.size();
    3.25  }
    3.26  
    3.27 +int Scene::get_num_kdnodes() const
    3.28 +{
    3.29 +	return kdtree_nodes(kdtree);
    3.30 +}
    3.31 +
    3.32  Material *Scene::get_materials()
    3.33  {
    3.34  	if(matlib.empty()) {
    3.35 @@ -159,6 +167,27 @@
    3.36  	return facebuf;
    3.37  }
    3.38  
    3.39 +static int find_node_index(const KDNode *node, const std::map<const KDNode*, int> &flatmap);
    3.40 +
    3.41 +static bool validate_flat_tree(const KDNode *node, const KDNodeGPU *kdbuf, int count, const std::map<const KDNode*, int> &flatmap)
    3.42 +{
    3.43 +	if(!node) return true;
    3.44 +
    3.45 +	int idx = find_node_index(node, flatmap);
    3.46 +	int left_idx = find_node_index(node->left, flatmap);
    3.47 +	int right_idx = find_node_index(node->right, flatmap);
    3.48 +
    3.49 +	if(kdbuf[idx].left != left_idx) {
    3.50 +		fprintf(stderr, "%d's left should be %d. it's: %d\n", idx, left_idx, kdbuf[idx].left);
    3.51 +		return false;
    3.52 +	}
    3.53 +	if(kdbuf[idx].right != right_idx) {
    3.54 +		fprintf(stderr, "%d's right should be %d. it's: %d\n", idx, right_idx, kdbuf[idx].right);
    3.55 +		return false;
    3.56 +	}
    3.57 +	return validate_flat_tree(node->left, kdbuf, count, flatmap) && validate_flat_tree(node->right, kdbuf, count, flatmap);
    3.58 +}
    3.59 +
    3.60  const KDNodeGPU *Scene::get_kdtree_buffer() const
    3.61  {
    3.62  	if(kdbuf) {
    3.63 @@ -169,29 +198,72 @@
    3.64  		((Scene*)this)->build_kdtree();
    3.65  	}
    3.66  
    3.67 -	int max_nodes = (int)pow(2, kdtree_depth(kdtree)) - 1;
    3.68 -	printf("allocating storage for the complete tree (%d)\n", max_nodes);
    3.69 +	/* we'll associate kdnodes with flattened array indices through this map as
    3.70 +	 * we add them so that the second child-index pass, will be able to find
    3.71 +	 * the children indices of each node.
    3.72 +	 */
    3.73 +	std::map<const KDNode*, int> flatmap;
    3.74 +	flatmap[0] = -1;
    3.75  
    3.76 -	kdbuf = new KDNodeGPU[max_nodes + 1];
    3.77 -	kdtree_gpu_flatten(kdbuf, 1, kdtree);
    3.78 +	int num_nodes = get_num_kdnodes();
    3.79 +	kdbuf = new KDNodeGPU[num_nodes];
    3.80 +
    3.81 +	int count = 0;
    3.82 +
    3.83 +	// first arrange the kdnodes into an array (flatten)
    3.84 +	flatten_kdtree(kdtree, kdbuf, &count, &flatmap);
    3.85 +
    3.86 +	// then fix all the left/right links to point to the correct nodes
    3.87 +	fix_child_links(kdtree, kdbuf, flatmap);
    3.88 +
    3.89 +	assert(validate_flat_tree(kdtree, kdbuf, count, flatmap));
    3.90 +
    3.91  	return kdbuf;
    3.92  }
    3.93  
    3.94 -static int ipow(int x, int n)
    3.95 +static void flatten_kdtree(const KDNode *node, KDNodeGPU *kdbuf, int *count, std::map<const KDNode*, int> *flatmap)
    3.96  {
    3.97 -	assert(n >= 0);
    3.98 +	int idx = (*count)++;
    3.99  
   3.100 -	int res = 1;
   3.101 -	for(int i=0; i<n; i++) {
   3.102 -		res *= x;
   3.103 +	// copy the node
   3.104 +	kdbuf[idx].aabb = node->aabb;
   3.105 +	for(size_t i=0; i<node->face_idx.size(); i++) {
   3.106 +		kdbuf[idx].face_idx[i] = node->face_idx[i];
   3.107  	}
   3.108 -	return res;
   3.109 +	kdbuf[idx].num_faces = node->face_idx.size();
   3.110 +
   3.111 +	// update the node* -> array-position mapping
   3.112 +	(*flatmap)[node] = idx;
   3.113 +
   3.114 +	// recurse to the left/right (if we're not in a leaf node)
   3.115 +	if(node->left) {
   3.116 +		assert(node->right);
   3.117 +
   3.118 +		flatten_kdtree(node->left, kdbuf, count, flatmap);
   3.119 +		flatten_kdtree(node->right, kdbuf, count, flatmap);
   3.120 +	}
   3.121  }
   3.122  
   3.123 -int Scene::get_kdtree_buffer_size() const
   3.124 +static int find_node_index(const KDNode *node, const std::map<const KDNode*, int> &flatmap)
   3.125  {
   3.126 -	// 2**depth - 1 nodes for the complete tree + 1 for the unused heap item 0.
   3.127 -	return ipow(2, kdtree_depth(kdtree)) * sizeof(KDNodeGPU);
   3.128 +	std::map<const KDNode*, int>::const_iterator it = flatmap.find(node);
   3.129 +	assert(it != flatmap.end());
   3.130 +	return it->second;
   3.131 +}
   3.132 +
   3.133 +static void fix_child_links(const KDNode *node, KDNodeGPU *kdbuf, const std::map<const KDNode*, int> &flatmap)
   3.134 +{
   3.135 +	if(!node) return;
   3.136 +
   3.137 +	int idx = find_node_index(node, flatmap);
   3.138 +
   3.139 +	kdbuf[idx].left = find_node_index(node->left, flatmap);
   3.140 +	if(!node->left && kdbuf[idx].left != -1) abort();
   3.141 +	kdbuf[idx].right = find_node_index(node->right, flatmap);
   3.142 +	if(!node->right && kdbuf[idx].right != -1) abort();
   3.143 +
   3.144 +	fix_child_links(node->left, kdbuf, flatmap);
   3.145 +	fix_child_links(node->right, kdbuf, flatmap);
   3.146  }
   3.147  
   3.148  void Scene::draw_kdtree() const
   3.149 @@ -442,33 +514,6 @@
   3.150  	return kdtree_nodes(node->left) + kdtree_nodes(node->right) + 1;
   3.151  }
   3.152  
   3.153 -#define MAX_FACES	(sizeof dest->face_idx / sizeof *dest->face_idx)
   3.154 -static void kdtree_gpu_flatten(KDNodeGPU *kdbuf, int idx, const KDNode *node)
   3.155 -{
   3.156 -	KDNodeGPU *dest = kdbuf + idx;
   3.157 -
   3.158 -	dest->aabb = node->aabb;
   3.159 -	dest->num_faces = 0;
   3.160 -
   3.161 -	for(size_t i=0; i<node->face_idx.size(); i++) {
   3.162 -		if(dest->num_faces >= (int)MAX_FACES) {
   3.163 -			fprintf(stderr, "kdtree_gpu_flatten WARNING: more than %d faces in node, skipping!\n", (int)MAX_FACES);
   3.164 -			break;
   3.165 -		}
   3.166 -		dest->face_idx[dest->num_faces++] = node->face_idx[i];
   3.167 -	}
   3.168 -
   3.169 -	if(node->left) {
   3.170 -		assert(node->right);
   3.171 -		assert(!dest->num_faces);
   3.172 -
   3.173 -		dest->num_faces = -1;
   3.174 -
   3.175 -		kdtree_gpu_flatten(kdbuf, idx * 2, node->left);
   3.176 -		kdtree_gpu_flatten(kdbuf, idx * 2 + 1, node->right);
   3.177 -	}
   3.178 -}
   3.179 -
   3.180  static void print_item_counts(const KDNode *node, int level)
   3.181  {
   3.182  	if(!node) return;

     4.1 --- a/src/scene.h	Tue Aug 24 05:47:04 2010 +0100
     4.2 +++ b/src/scene.h	Thu Aug 26 20:24:07 2010 +0100
     4.3 @@ -54,7 +54,8 @@
     4.4  	AABBox aabb;
     4.5  	int face_idx[32];
     4.6  	int num_faces;
     4.7 -	int padding[3];
     4.8 +	int left, right;
     4.9 +	int padding;
    4.10  };
    4.11  
    4.12  
    4.13 @@ -75,8 +76,9 @@
    4.14  
    4.15  	bool add_mesh(Mesh *m);
    4.16  	int get_num_meshes() const;
    4.17 +	int get_num_faces() const;
    4.18  	int get_num_materials() const;
    4.19 -	int get_num_faces() const;
    4.20 +	int get_num_kdnodes() const;
    4.21  
    4.22  	Material *get_materials();
    4.23  	const Material *get_materials() const;
    4.24 @@ -86,7 +88,6 @@
    4.25  
    4.26  	const Face *get_face_buffer() const;
    4.27  	const KDNodeGPU *get_kdtree_buffer() const;
    4.28 -	int get_kdtree_buffer_size() const;
    4.29  
    4.30  	void draw_kdtree() const;
    4.31  	bool build_kdtree();