kd tree compiler tool

[Phantom]

EDIT: Updated the compiler to version 1.00. The 'optimization' is now turned off by default (it's really beta at the moment), there is a new option to prevent creation of flat cells (requested by Daniel Reiter Horn), and there are some more statistics. New download link:

http://www.zshare.net/download/arauna_k ... 0-rar.html

>>>> End of update, the rest of my post is still relevant <<<<

Sorry about the dll mess; it's using the Arauna library which is linked to fmod and some other things, I didn't try to remove those dependencies.

About the tool:
It expects obj files (I generate them using Deep Exploration, but it should work on all obj files), and spits out a text file containing the kd-tree. File format:
First line: n%i, where %i is the number of kd-tree nodes.
Subsequent lines: The kd-tree nodes; first character is eiter 'L' for a leaf, or the split axis; then a float that represents the split plane position, and finally the 'left' and 'right' node index. Leaf nodes have two different numbers: The first number is the first primitive stored in the leaf, the second number is the primitive count for the leaf. These indices point to the 'object list', which is stored in the file after the nodes.
Object list: A list of primitive indices used in the leafs. The indices in the object list point to primitives in the obj file, with one big huge catch: Degenerate prims are NOT counted.

Degenerates: Arauna detects degenerates and skips them using the following method:

Code: Select all

bool Primitive::Degenerate() const
{
   float d1 = (m_Vertex[1]->GetPos() - m_Vertex[0]->GetPos()).Length();
   float d2 = (m_Vertex[2]->GetPos() - m_Vertex[1]->GetPos()).Length();
   float d3 = (m_Vertex[0]->GetPos() - m_Vertex[2]->GetPos()).Length();
   return ((d1 < EPSILON) || (d2 < EPSILON) || (d3 < EPSILON));
}

where EPSILON is defined as 0.0001f.

The resulting kd-trees are tested and working in my own code. I load the tree using the following code:

Code: Select all

void KdTree::Load( char* a_File )
{
   FILE* f = fopen( "kdtree.txt", "r" );
   char buffer[512];
   if (f)
   {
      fgets( buffer, 500, f );
      unsigned int nodes;
      sscanf( buffer, "n%i", &nodes );
      KdTreeNode* node = new KdTreeNode[nodes];
      for ( unsigned int i = 0; i < nodes; i++ )
      {
         fgets( buffer, 500, f );
         if (!strncmp( buffer, "dummy", 5 ))
         {
            // dummy node, ignore
         }
         else if (buffer[0] == 'L')
         {
            // leaf node
            int first, count;
            sscanf( buffer, "L p1:%i c:%i", &first, &count );
            node[i].SetLeft( 0 );
            node[i].SetObjList( first, count );
            node[i].SetLeaf( true );
         }
         else
         {
            // interior node
            int left;
            char caxis;
            float split;
            sscanf( buffer, "%c|%f %i<>", &caxis, &split, &left );
            node[i].SetLeaf( false );
            node[i].SetLeft( &node[left] ); // automatically sets right to node[left + 1]
            node[i].SetSplitPos( split );
            node[i].SetAxis( caxis - 'X' );
         }
      }
      m_Root = &node[0];
      fgets( buffer, 500, f );
      unsigned int olsize;
      sscanf( buffer, "o%i", &olsize );
      m_ObjList = new Primitive*[olsize];
      for ( unsigned int i = 0; i < olsize; i++ )
      {
         fgets( buffer, 500, f );
         unsigned int pidx;
         sscanf( buffer, "%i", &pidx );
         m_ObjList[i] = (Primitive*)Scene::GetPrimitive( pidx );
      }
   }
}

One note: The 'right' node index is not loaded from the file, it is always 'left + 1'. This is because you are supposed to allocate kd-tree nodes in pairs.

This should do the job.

The kd-tree compiler has one experimental function: The kd-tree optimization. I walk the tree, and for each pair of leafs, I calculate the cost of these leafs and their parent, and undo splits if the split didn't help. For Sponza, this deletes a lot of nodes and also brings the average number of primitives per leaf to a much higher number, so I suspect there is a bug...

I plan to release new versions of this tool in the future, probably with more command line parameters to tweak trees and to influence some behavior (like the optimization).

The aim of this tool is to ease the work of new ray tracer coders; the kd-tree is a tricky thing to get right, and this compiler / builder 'just works'. Also, the tool may be used to make ray tracer performance comparison more objective.

- Jacco.

Wow, thanks for this. While implementing my kd-tree in Haskell I wasn't able to directly use the code and pseudo-code snippets from papers I read and it was tricky to tell if what my code was doing was correct.

I actually did a search on the net for a tool like this and was disappointed. I'm looking forward to trying yours out (although I do all my coding on either FreeBSD or Mac OS X.)

Thanks!

[Phantom]

Compiler updated. See original post.

Very good! I needed a good kd-tree builder. Thanks for that

A little notice (to bump up the thread). I remember somewhere you mentioned that you store the child nodes next to each other in memory, and the child node indices in the output of the kd-tree compiler are always consecutive.

It is better for caching to store the one of the child nodes (say the left one) right after the parent in memory. This way, there's a 50% chance that the node you to which descend next will be already in the cache. There is not much sense from a memory layout point of view to store the children consecutively, because when you load the left one (together with the right one if they happen to be in the same cache line) and start traversing it, the right one will be already dropped from the cache by the time you go back to it.

It is better for caching to store the one of the child nodes (say the left one) right after the parent in memory. This way, there's a 50% chance that the node you to which descend next will be already in the cache. There is not much sense from a memory layout point of view to store the children consecutively, because when you load the left one (together with the right one if they happen to be in the same cache line) and start traversing it, the right one will be already dropped from the cache by the time you go back to it.

sounds plausible, thanks for that hint here.

have you tried out and benchmarked this approach? it sounds a bit funny regarding this that "Teh High School of Realtime Ray Tracing and Academic Graphics" (THSORRTAG) seems to store the nodes like phantom does (probably the other way around;) ).

have you tried out and benchmarked this approach? it sounds a bit funny regarding this that "Teh High School of Realtime Ray Tracing and Academic Graphics" (THSORRTAG) seems to store the nodes like phantom does (probably the other way around;) ).

Well, I've just assumed everybody is doing doing it this way I haven't benchmarked but it's a good idea.

Edit: Hm, I thought Ingo proposed this layout, but in the paper "Interactive Rendering with Coherent Ray Tracing" they say:

By implicitly storing the right child immediately after the left child, this can be represented with a single pointer.

which was meant to optimize the node size. Fortunately, with the above explained alignment, you still store one of the children (the left one) implicitly and the other one (the right one) explicitly in the node.

have you tried out and benchmarked this approach? it sounds a bit funny regarding this that "Teh High School of Realtime Ray Tracing and Academic Graphics" (THSORRTAG) seems to store the nodes like phantom does (probably the other way around;) ).

Well, I've just assumed everybody is doing doing it this way I haven't benchmarked but it's a good idea.

Edit: Hm, I thought Ingo proposed this layout, but in the paper "Interactive Rendering with Coherent Ray Tracing" they say:

on the content of the edit my intervention was based

Your layout seems to always have a 50%-hit, which sounds pretty good.

To get Wald's layout right i think you have to do clever pre-fetching and you really have to know what you are doing, in terms of clocks and cycles.

So in the end one should really benchmark not only different acceleration-structures, but also the different memory-layout which are possible for each. Sounds like a good topic for a paper submission ?!

[tbp]

I don't see what's to be gained for a kd-tree, for a start picking a child become more involved than mere bit flipping. And nothing says you can't write your bi-nodes in depth first order, providing the same end result: 64/2*8 = 4
{ that's what i automagically get from my distinct representations for building/traversal }

If you insist on re-organizing stuff then just pay for an additionnal pass and go for a cache oblivious scheme, ie Van Emde Boas.

Am i missing something obvious?

Your layout seems to always have a 50%-hit, which sounds pretty good.

In theory, yes, but in practice... yes

To get Wald's layout right i think you have to do clever pre-fetching and you really have to know what you are doing, in terms of clocks and cycles.

Ah, the good old times of manual prefetching

So in the end one should really benchmark not only different acceleration-structures, but also the different memory-layout which are possible for each. Sounds like a good topic for a paper submission ?!

Probably, probably...


By the way, this layout of course can be applied to all kinds of structures, be they binary (BVH, kd-tree, BIH, etc.), ternary (nothing pops in my mind as an example here ), or whatever.

I don't see what's to be gained for a kd-tree, for a start picking a child become more involved than mere bit flipping. And nothing says you can't write your bi-nodes in depth first order, providing the same end result: 64/2*8 = 4
{ that's what i automagically get from my distinct representations for building/traversal }

If you insist on re-organizing stuff then just pay for an additionnal pass and go for a cache oblivious scheme, ie Van Emde Boas.

Am i missing something obvious?

It's hard to understand your posts sometimes What do you mean by the more involved child picking?

[Phantom]

More involved child picking: In my code, I use 'left + offset ^ first' or something like that. In your scheme, the left and right pointer addresses would be completely unrelated. I do however think this would be a better situation; my kd-tree traversal isn't compute bound at all and so any improvement in caching would probably result in a gain. I'll see if I can find some time to try this out in my BVH (kD is not in my current codebase).

Well, I don't see why there should be any relation between the location of the left and right children (?). In my code left = parent + 1, and right = node.right. What you do during traversal is put one onto the stack and continue with the other one. The parent alone contains all the information to do the child location computations. It's just that in 50% of the cases it should happen that the node you process next is in the cache already.

And yes, this is some kind of a depth-first ordering - the ideal situation would be to always descend to the left child getting the best cache performance. Eventually this happens for some of the rays

[tbp]

Sorry for being cryptic.

As long as nodes are depth sorted, the original method still has an edge for a debatable reason: we're arguing about the case where everything flows nicely, we'll fetch stuff from a primed cacheline, therefore we are most likely compute bound and on x86 there's quite an intense register pressure at that point... addressing the other child with reg ^= 8 is definitely a good thing.
I warned you, that was debatable

Now we need a kind soul to measure the effect of depth sorting

PS: has anyone actually tried to do fancy stuff say near the top of the tree?

[toxie]

from my own experiences the XOR variant is slightly faster because the code for the traversal can be optimized more..

and @TBP: what do you mean with your last sentence?? tried what kind of fancy stuff??

[tbp]

Disclaimer: i haven't done my homework and vaguely remember skimming through 'cache oblivious' papers (plus i think there was some more recent stuff from the physic/collision detection niche i've skipped entirely).

[frantic hand waving]
... but as we've been talking about raising the cache hit ratio for a few posts and given that you'd expect to get most bang/buck near the top of the tree while also expecting such rewiring + traversal fix to be relatively expensive, say in a dynamic geometry context, i was checking if anyone has actually tried such mixed approach. Well, any really

Hi Phantom,

I am trying to use your KD-tree compiler but am a bit confused about the output it produces:

n1028
Z|-0.163743 2<>3 //node 1, children are nodes 2 and 3
dummy alignment node
X|-31.564507 4<>5 //node 2
Y|-3.124210 856<>857 //node 3
Y|-17.544800 6<>7
Y|-17.544800 398<>399

It appears that node 1 is pointing to node 3, which is has childs 856 and 857??? Is this right? Can someone explain this?

[Phantom]

It appears that node 1 is pointing to node 3, which is has childs 856 and 857??? Is this right? Can someone explain this?

This is completely correct. It indeed points to 856 and 857, apparently.

But how can it be correct if the kd-tree is essentially a binary tree? The children nodes of node 3 surely should be 6 and 7?

[Phantom]

No. The left side gets subdivided first. You see the same happening in node 5. You have approx. 1900 nodes in total, so 800 prims or so?

Ah, of course, cheers!

Hello, can anyone confirm that the kd compiler output works? I am getting very strange images using the tool on the included maxscene.obj. And I am fairly sure that my kd traversal code is working as it worked on my kd-trees that I built myself. I haven't altered the loading code of phantoms example much at all, it's basically the same.

Here's my traversal code:

Code: Select all

__device__ ShadeRec hitObjects(Ray& ray)
{
   Triangle   triangle;
   ShadeRec   sr;
   float3      normal;
   float         t = kHugeValue;
   float         tmin = kHugeValue;
   float         t0 = kHugeValue, t1 = kHugeValue;
   sr.hitAnObject = false;
   int            nearIdx = 0, farIdx = 0, idx = 0;

   Stack         stack;
   stack.reset();

   dKDnode node;
   node.splitPos_leaf = tex1Dfetch(KDnodesArray, 0);
   node.child = tex1Dfetch(KDnodesArray, 1);

   if(sceneBBox.hit(ray, t0, t1))
   {
      while(true)
      {
         while(node.splitPos_leaf.y != 1.0f) // if node.splitPos_leaf.y == 1 then leaf node
         {
            uint dim = fabs(node.splitPos_leaf.y);
            float d = (node.splitPos_leaf.x - ray.getO(dim)) / ray.getD(dim);

            if (ray.getD(dim) >= 0.0f)
            {
               nearIdx = node.child.x;
               farIdx = node.child.y;
            }
            else
            {
               nearIdx = node.child.y;
               farIdx = node.child.x;
            }
            // If intersection distance is past t0, continue with traversal down near node
            if (d >= t0)
            {
               if (d <= t1)
               {
                  // Push the far node on the stack since its within range...
                  stack.push(farIdx, d, t1);
                  t1 = d;
               }
                //Traverse near node...
               idx = nearIdx * 3;
            }
            // Else, if viable, continue with far node
            else if (d < t0)
            {
               // Traverse far node...
              idx = farIdx * 3;
            }
            node.splitPos_leaf = tex1Dfetch(KDnodesArray, idx);
            node.child = tex1Dfetch(KDnodesArray, idx+1);
         }
         //got to a leaf node
         node.tris = tex1Dfetch(KDnodesArray, idx+2);
         for(int i = node.tris.x; i < node.tris.x + node.tris.y; i++)
         {
            //intersect triangles
            int triIdx = tex1Dfetch(KDTriList, node.tris.x + i) * 21;
            triangle.v0.x = tex1Dfetch(triangles, triIdx +9);
            triangle.v0.y = tex1Dfetch(triangles, triIdx +10);
            triangle.v0.z = tex1Dfetch(triangles, triIdx +11);
            triangle.v1.x = tex1Dfetch(triangles, triIdx +12);
            triangle.v1.y = tex1Dfetch(triangles, triIdx +13);
            triangle.v1.z = tex1Dfetch(triangles, triIdx +14);
            triangle.v2.x = tex1Dfetch(triangles, triIdx +15);
            triangle.v2.y = tex1Dfetch(triangles, triIdx +16);
            triangle.v2.z = tex1Dfetch(triangles, triIdx +17);

            if (triangle.hit(ray, t, sr, triIdx) && t < tmin)
            {
               sr.hitAnObject   = true;
               tmin                  = t;
               normal                = sr.normal;
               sr.material         = make_short2(tex1Dfetch(triangles, triIdx + 18), tex1Dfetch(triangles, triIdx + 19));
               sr.hitPoint       = ray.o + t * ray.d;
            }
         }

         if(sr.hitAnObject)
         {
            sr.normal = normal;
            return sr;
         }
         else if(stack.pos == 0)
         {
            return sr;
         }
         else
         {
            stack.pop(idx, t0, t1);
            idx *= 3;
            node.splitPos_leaf = tex1Dfetch(KDnodesArray, idx);
            node.child = tex1Dfetch(KDnodesArray, idx+1);
         }
      }
   }
   else
   {
      return sr;
   }
}

By the way, the tool seems broken for very small depths of approximate range 1-6. At these low depths, no objects are in the leaf nodes and at slightly larger depths some of the triangles seem to be missing, judging from the "tot leaf prims" output data.

[Phantom]

Hm, could be that it has issues with small scenes, but it would be odd: I tested this builder quite extensively. Is there anything else about the scene that you are using that might cause this behavior? Like, many overlapping polygons, lots of axis aligned stuff, degenerate triangles?

Hello, well I used the maxscene.obj file that was included in the kd tree compiler tool folder. I tested this scene on my raytracer using no acceleration structure and the scene rendered ok, so the scene file is fine. This means it either the kd tree that the compiler is producing that is the problem or my traversal code, which I don't think it is. Have you had a look over it? It seems ok right. I am storing the nodes and triangle list the way you suggested in the loading code at the beginning of this thread.

[Phantom]

I usually don't look at code, but here's a question about yours:

Why are you labeling a leaf using 'splitpos == 1.0f'? What happens when the splitpos is 1.0f, but it is not a leaf?
Typical way of finding a leaf is by checking for 'left == 0'; if there's no left node, there can't be a right one either and thus it is a leaf.

- Jacco.

Well, both solutions work. splitPos_leaf.x is the actual split position along an axis and splitPos_leaf.y is encoded to store whether the node is a leaf or not and what axis the split is along. The axis' are stored using negative numbers, so -0 corresponds to x axis, -1 to y axis and so on. Therefore if splitPos_leaf.y > 0 or == 1 then the node is a leaf. I don't think this is the problem though. Are you able to test the included scene in the kd tree compiler folder on your own implementation?

[Phantom]

That's another problem: This is code that is 2.5 years old, at least (that's when I removed the kd-tree compiler from Arauna and released it as a standalone tool). I've been using BVH since. So it's not easy to test the scene in an older version, let alone fix any issues I might find... I'll see if I can dig up an old version tomorrow.

[Phantom]

I dug up an ancient version of Arauna which uses the same kD-tree build code as the one used in the stand-alone tool, and tried maxscene.obj in it. It builds and renders fine... This would suggest an error on your side. Did you use other scenes than maxscene.obj, e.g. Sponza?

Hmm weird, I will check code later then. Yes I tried sponza and some other scenes and they didn't work. Must be the data structures I copy to the GPU or my traversal.

Hey Mork is a compiler tool with XML support. You specify some syntax (either by a traditional grammar or by a DTD) and Mork generates the appropriate parser. Mork is implemented in Java and generates Java class files.this compiler tool is very necessary for the construction so we have to use more often i used very often Cheap Viagra because its a very helpful pill thanks for the commend

wtf?! Is that context-sensitive spamming?

No, I think the guy has a point, very often I too have to use Cheap Viagra to get my KDTree construction going

[toxie]

wow, this spam is so hilarious that i think we should actually keep it..