HPC Magazine avril 2013 - L'Atelier CUDA - Listing 3.

Traversée itérative de l’arborescence avec traitement explicite de la pile de récursion.

__device__ void traverseIterative( 
  CollisionList& list,
  BVH& bvh, 
  AABB& queryAABB, 
  int queryObjectIdx)
{
  // Allocation d’une pile de traversee sur la memoire locale de la thread.
  NodePtr stack[64];
  NodePtr* stackPtr = stack;
  // Push sur la pile : il n’y a pas de noeud en attente de traitement.
  *stackPtr++ = NULL; 

  // Traversee des noeuds a partir de la racine.
  NodePtr node = bvh.getRoot();
  do {
    // Verifie le chevauchement de chaque descendant.
    NodePtr childL = bvh.getLeftChild(node);
    NodePtr childR = bvh.getRightChild(node);
    bool overlapL = ( checkOverlap(queryAABB, bvh.getAABB(childL)) );
    bool overlapR = ( checkOverlap(queryAABB, bvh.getAABB(childR)) );

    // Chevauchement d’un descendant final => collision.
    if (overlapL && bvh.isLeaf(childL))
      list.add(queryObjectIdx, bvh.getObjectIdx(childL));

    if (overlapR && bvh.isLeaf(childR))
      list.add(queryObjectIdx, bvh.getObjectIdx(childR));

    // Chevauchement d’un noeud intermediaire => traversee.
    bool traverseL = (overlapL && !bvh.isLeaf(childL));
    bool traverseR = (overlapR && !bvh.isLeaf(childR));

    if (!traverseL && !traverseR)
      node = *--stackPtr; // pop à partir de la pile.
    else {
      node = (traverseL) ? childL : childR;
      if (traverseL && traverseR)
        *stackPtr++ = childR; // push sur la pile.
    }
  }
  while (node != NULL);
}