1.1 --- a/rt.cl	Sat Aug 21 23:57:19 2010 +0100
     1.2 +++ b/rt.cl	Sun Aug 22 00:50:47 2010 +0100
     1.3 @@ -170,34 +170,40 @@
     1.4  	sp0.obj = 0;
     1.5  
     1.6  	int idxstack[STACK_SIZE];
     1.7 -	int sp = 0;			// points at the topmost element of the stack
     1.8 -	idxstack[sp] = 1;	// root at tree[1] (heap)
     1.9 +	int top = 0;			// points after the topmost element of the stack
    1.10 +	idxstack[top++] = 1;	// root at tree[1] (heap)
    1.11  
    1.12 -	printf("check intersection\n");
    1.13 +	while(top > 0) {
    1.14 +		int idx = idxstack[--top];	// remove this index from the stack and process it
    1.15  
    1.16 -	while(sp >= 0) {
    1.17 -		int idx = idxstack[sp--];	// remove this index from the stack and process it
    1.18 +		global const struct KDNode *node = scn->kdtree + idx;
    1.19  
    1.20 -		global struct KDNode *node = scn->kdtree + idx;
    1.21 -		printf("idx: %d (%p) num_faces: %d\n", idx, node, node->num_faces);
    1.22 +		/*if(get_global_id(0) == 0) {
    1.23 +			for(int i=0; i<top+1; i++) {
    1.24 +				printf("   ");
    1.25 +			}
    1.26 +			printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);
    1.27 +		}*/
    1.28  
    1.29  		if(intersect_aabb(ray, node->aabb)) {
    1.30 -			// leaf node ...
    1.31  			if(node->num_faces >= 0) {
    1.32 -				// check each face in turn and update the nearest intersection as needed
    1.33 +				// leaf node... check each face in turn and update the nearest intersection as needed
    1.34  				for(int i=0; i<node->num_faces; i++) {
    1.35 -					struct SurfPoint sp;
    1.36 +					struct SurfPoint spt;
    1.37  					int fidx = node->face_idx[i];
    1.38  
    1.39 -					if(intersect(ray, scn->faces + fidx, &sp) && sp.t < sp0.t) {
    1.40 -						sp0 = sp;
    1.41 +					if(intersect(ray, scn->faces + fidx, &spt) && spt.t < sp0.t) {
    1.42 +						sp0 = spt;
    1.43  					}
    1.44  				}
    1.45 +			} else {
    1.46 +				// internal node... recurse to the children
    1.47 +				/*if(get_global_id(0) == 0) {
    1.48 +					printf("pushing %d's children %d and %d\n", idx, idx * 2, idx * 2 + 1);
    1.49 +				}*/
    1.50 +				idxstack[top++] = idx * 2;
    1.51 +				idxstack[top++] = idx * 2 + 1;
    1.52  			}
    1.53 -		} else {
    1.54 -			// internal node ... recurse to the children
    1.55 -			idxstack[++sp] = idx * 2;
    1.56 -			idxstack[++sp] = idx * 2 + 1;
    1.57  		}
    1.58  	}
    1.59  

     2.1 --- a/src/scene.cc	Sat Aug 21 23:57:19 2010 +0100
     2.2 +++ b/src/scene.cc	Sun Aug 22 00:50:47 2010 +0100
     2.3 @@ -472,10 +472,12 @@
     2.4  
     2.5  	if(node->left) {
     2.6  		assert(node->right);
     2.7 +		assert(!dest->num_faces);
     2.8 +
     2.9 +		dest->num_faces = -1;
    2.10 +
    2.11  		kdtree_gpu_flatten(kdbuf, idx * 2, node->left, facebuf);
    2.12  		kdtree_gpu_flatten(kdbuf, idx * 2 + 1, node->right, facebuf);
    2.13 -	} else {
    2.14 -		dest->num_faces = -1;
    2.15  	}
    2.16  }
    2.17