# HG changeset patch # User John Tsiombikas # Date 1282850647 -3600 # Node ID 7d77ded5f8906fce77503371d5003d2fe9d929ca # Parent 931d13b72f838c50738d66a204d338a1dfede35a stopped using a heap to flatten the kdtree. added explicit left/right indices diff -r 931d13b72f83 -r 7d77ded5f890 rt.cl --- a/rt.cl Tue Aug 24 05:47:04 2010 +0100 +++ b/rt.cl Thu Aug 26 20:24:07 2010 +0100 @@ -62,7 +62,8 @@ struct AABBox aabb; int face_idx[32]; int num_faces; - int padding[3]; + int left, right; + int padding; }; #define MIN_ENERGY 0.001 @@ -117,9 +118,9 @@ ray.origin = sp.pos; ray.dir = reflect(-ray.dir, sp.norm); - energy *= sp.mat.ks * sp.mat.kr; + energy *= refl_col; } else { - iter = INT_MAX - 1; // to break out of the loop + break; } } @@ -171,7 +172,7 @@ int idxstack[STACK_SIZE]; int top = 0; // points after the topmost element of the stack - idxstack[top++] = 1; // root at tree[1] (heap) + idxstack[top++] = 0; // root at tree[0] while(top > 0) { int idx = idxstack[--top]; // remove this index from the stack and process it @@ -186,7 +187,7 @@ }*/ if(intersect_aabb(ray, node->aabb)) { - if(node->num_faces >= 0) { + if(node->left == -1) { // leaf node... check each face in turn and update the nearest intersection as needed for(int i=0; inum_faces; i++) { struct SurfPoint spt; @@ -199,10 +200,10 @@ } else { // internal node... recurse to the children /*if(get_global_id(0) == 0) { - printf("pushing %d's children %d and %d\n", idx, idx * 2, idx * 2 + 1); + printf("pushing %d's children %d and %d\n", idx, node->left, node->right); }*/ - idxstack[top++] = idx * 2; - idxstack[top++] = idx * 2 + 1; + idxstack[top++] = node->left; + idxstack[top++] = node->right; } } } diff -r 931d13b72f83 -r 7d77ded5f890 src/rt.cc --- a/src/rt.cc Tue Aug 24 05:47:04 2010 +0100 +++ b/src/rt.cc Thu Aug 26 20:24:07 2010 +0100 @@ -102,7 +102,7 @@ prog->set_arg_buffer(KARG_PRIM_RAYS, ARG_RD, xsz * ysz * sizeof *prim_rays, prim_rays); prog->set_arg_buffer(KARG_XFORM, ARG_RD, 16 * sizeof(float)); prog->set_arg_buffer(KARG_INVTRANS_XFORM, ARG_RD, 16 * sizeof(float)); - prog->set_arg_buffer(KARG_KDTREE, ARG_RD, scn->get_kdtree_buffer_size(), kdbuf); + prog->set_arg_buffer(KARG_KDTREE, ARG_RD, scn->get_num_kdnodes() * sizeof *kdbuf, kdbuf); if(prog->get_num_args() < NUM_KERNEL_ARGS) { return false; diff -r 931d13b72f83 -r 7d77ded5f890 src/scene.cc --- a/src/scene.cc Tue Aug 24 05:47:04 2010 +0100 +++ b/src/scene.cc Thu Aug 26 20:24:07 2010 +0100 @@ -1,15 +1,18 @@ +#include #include #include #include +#include #include "scene.h" #include "ogl.h" +static void flatten_kdtree(const KDNode *node, KDNodeGPU *kdbuf, int *count, std::map *flatmap); +static void fix_child_links(const KDNode *node, KDNodeGPU *kdbuf, const std::map &flatmap); static void draw_kdtree(const KDNode *node, int level = 0); static bool build_kdtree(KDNode *kd, const Face *faces, int level = 0); static float eval_cost(const Face *faces, const int *face_idx, int num_faces, const AABBox &aabb, int axis); static void free_kdtree(KDNode *node); -static void kdtree_gpu_flatten(KDNodeGPU *kdbuf, int idx, const KDNode *node); static void print_item_counts(const KDNode *node, int level); @@ -123,6 +126,11 @@ return (int)matlib.size(); } +int Scene::get_num_kdnodes() const +{ + return kdtree_nodes(kdtree); +} + Material *Scene::get_materials() { if(matlib.empty()) { @@ -159,6 +167,27 @@ return facebuf; } +static int find_node_index(const KDNode *node, const std::map &flatmap); + +static bool validate_flat_tree(const KDNode *node, const KDNodeGPU *kdbuf, int count, const std::map &flatmap) +{ + if(!node) return true; + + int idx = find_node_index(node, flatmap); + int left_idx = find_node_index(node->left, flatmap); + int right_idx = find_node_index(node->right, flatmap); + + if(kdbuf[idx].left != left_idx) { + fprintf(stderr, "%d's left should be %d. it's: %d\n", idx, left_idx, kdbuf[idx].left); + return false; + } + if(kdbuf[idx].right != right_idx) { + fprintf(stderr, "%d's right should be %d. it's: %d\n", idx, right_idx, kdbuf[idx].right); + return false; + } + return validate_flat_tree(node->left, kdbuf, count, flatmap) && validate_flat_tree(node->right, kdbuf, count, flatmap); +} + const KDNodeGPU *Scene::get_kdtree_buffer() const { if(kdbuf) { @@ -169,29 +198,72 @@ ((Scene*)this)->build_kdtree(); } - int max_nodes = (int)pow(2, kdtree_depth(kdtree)) - 1; - printf("allocating storage for the complete tree (%d)\n", max_nodes); + /* we'll associate kdnodes with flattened array indices through this map as + * we add them so that the second child-index pass, will be able to find + * the children indices of each node. + */ + std::map flatmap; + flatmap[0] = -1; - kdbuf = new KDNodeGPU[max_nodes + 1]; - kdtree_gpu_flatten(kdbuf, 1, kdtree); + int num_nodes = get_num_kdnodes(); + kdbuf = new KDNodeGPU[num_nodes]; + + int count = 0; + + // first arrange the kdnodes into an array (flatten) + flatten_kdtree(kdtree, kdbuf, &count, &flatmap); + + // then fix all the left/right links to point to the correct nodes + fix_child_links(kdtree, kdbuf, flatmap); + + assert(validate_flat_tree(kdtree, kdbuf, count, flatmap)); + return kdbuf; } -static int ipow(int x, int n) +static void flatten_kdtree(const KDNode *node, KDNodeGPU *kdbuf, int *count, std::map *flatmap) { - assert(n >= 0); + int idx = (*count)++; - int res = 1; - for(int i=0; iaabb; + for(size_t i=0; iface_idx.size(); i++) { + kdbuf[idx].face_idx[i] = node->face_idx[i]; } - return res; + kdbuf[idx].num_faces = node->face_idx.size(); + + // update the node* -> array-position mapping + (*flatmap)[node] = idx; + + // recurse to the left/right (if we're not in a leaf node) + if(node->left) { + assert(node->right); + + flatten_kdtree(node->left, kdbuf, count, flatmap); + flatten_kdtree(node->right, kdbuf, count, flatmap); + } } -int Scene::get_kdtree_buffer_size() const +static int find_node_index(const KDNode *node, const std::map &flatmap) { - // 2**depth - 1 nodes for the complete tree + 1 for the unused heap item 0. - return ipow(2, kdtree_depth(kdtree)) * sizeof(KDNodeGPU); + std::map::const_iterator it = flatmap.find(node); + assert(it != flatmap.end()); + return it->second; +} + +static void fix_child_links(const KDNode *node, KDNodeGPU *kdbuf, const std::map &flatmap) +{ + if(!node) return; + + int idx = find_node_index(node, flatmap); + + kdbuf[idx].left = find_node_index(node->left, flatmap); + if(!node->left && kdbuf[idx].left != -1) abort(); + kdbuf[idx].right = find_node_index(node->right, flatmap); + if(!node->right && kdbuf[idx].right != -1) abort(); + + fix_child_links(node->left, kdbuf, flatmap); + fix_child_links(node->right, kdbuf, flatmap); } void Scene::draw_kdtree() const @@ -442,33 +514,6 @@ return kdtree_nodes(node->left) + kdtree_nodes(node->right) + 1; } -#define MAX_FACES (sizeof dest->face_idx / sizeof *dest->face_idx) -static void kdtree_gpu_flatten(KDNodeGPU *kdbuf, int idx, const KDNode *node) -{ - KDNodeGPU *dest = kdbuf + idx; - - dest->aabb = node->aabb; - dest->num_faces = 0; - - for(size_t i=0; iface_idx.size(); i++) { - if(dest->num_faces >= (int)MAX_FACES) { - fprintf(stderr, "kdtree_gpu_flatten WARNING: more than %d faces in node, skipping!\n", (int)MAX_FACES); - break; - } - dest->face_idx[dest->num_faces++] = node->face_idx[i]; - } - - if(node->left) { - assert(node->right); - assert(!dest->num_faces); - - dest->num_faces = -1; - - kdtree_gpu_flatten(kdbuf, idx * 2, node->left); - kdtree_gpu_flatten(kdbuf, idx * 2 + 1, node->right); - } -} - static void print_item_counts(const KDNode *node, int level) { if(!node) return; diff -r 931d13b72f83 -r 7d77ded5f890 src/scene.h --- a/src/scene.h Tue Aug 24 05:47:04 2010 +0100 +++ b/src/scene.h Thu Aug 26 20:24:07 2010 +0100 @@ -54,7 +54,8 @@ AABBox aabb; int face_idx[32]; int num_faces; - int padding[3]; + int left, right; + int padding; }; @@ -75,8 +76,9 @@ bool add_mesh(Mesh *m); int get_num_meshes() const; + int get_num_faces() const; int get_num_materials() const; - int get_num_faces() const; + int get_num_kdnodes() const; Material *get_materials(); const Material *get_materials() const; @@ -86,7 +88,6 @@ const Face *get_face_buffer() const; const KDNodeGPU *get_kdtree_buffer() const; - int get_kdtree_buffer_size() const; void draw_kdtree() const; bool build_kdtree();