BVH Acceleration

Bounding Volume Hierarchy (BVH) is one of the most critical acceleration structures in ray tracing.

Overview

BVH organizes scene objects into a hierarchical bounding box tree, reducing ray intersection complexity from O(N) to O(log N).

SAH Split Strategy

Surface Area Heuristic (SAH) is a cost function for selecting optimal split positions:

$$C = C_{trav} + \frac{SA(L)}{SA(N)} \cdot |L| \cdot C_{isect} + \frac{SA(R)}{SA(N)} \cdot |R| \cdot C_{isect}$$

Where:

• $SA$ represents the bounding box surface area
• $|L|$, $|R|$ are the primitive counts in left/right subtrees
• $C_{trav}$ is the node traversal cost
• $C_{isect}$ is the intersection test cost

Build Algorithm

BVH Build (SAH)

Input: Primitive set P Output: BVH root node

1. if |P| ≤ leaf threshold then
2. return create leaf node(P)
3. end if
4. Choose split axis (x/y/z) minimizing SAH cost
5. Partition P into P_left, P_right
6. node.left ← BVH-Build(P_left)
7. node.right ← BVH-Build(P_right)
8. return node

Traversal Algorithm

BVH Traversal (Stack-based)

Input: Ray R, BVH root node Output: Closest intersection hit

1. stack.push(root)
2. while stack not empty do
3. node ← stack.pop()
4. if R does not intersect node.aabb then continue
5. if node is leaf then

6. Test R against all primitives in node
   

7. else

8. stack.push(node.left)
   

9. stack.push(node.right)
   

10. end if
11. end while

Performance Analysis

Scene	Spheres	Brute Force	BVH	Speedup
Demo	10	10	4	2.5×
Random	100	100	7	14×
Complex	1000	1000	10	100×

Implementation

// BVH Node Structure
struct BVHNode {
    AABB bounds;
    int left;    // Left child index (negative indicates leaf)
    int right;   // Right child index
    int start;   // Leaf node: primitive start index
    int count;   // Leaf node: primitive count
};

// BVH Traversal Kernel
__device__ bool traverse_bvh(
    const Ray& ray,
    const BVHNode* nodes,
    const Sphere* spheres,
    HitRecord& closest_hit
) {
    int stack[64];
    int stack_ptr = 0;
    stack[stack_ptr++] = 0;  // Root node

    while (stack_ptr > 0) {
        int node_idx = stack[--stack_ptr];
        const BVHNode& node = nodes[node_idx];

        if (!intersect_aabb(ray, node.bounds)) continue;

        if (node.left < 0) {  // Leaf node
            for (int i = 0; i < node.count; ++i) {
                // Test primitive intersection
            }
        } else {
            stack[stack_ptr++] = node.left;
            stack[stack_ptr++] = node.right;
        }
    }
    return closest_hit.valid;
}

Current Implementation Details

The current BVH implementation uses:

• Midpoint Split: Objects sorted along split axis, divided at midpoint
• CPU Construction: BVH built on host, copied to device
• Stack-based Traversal: GPU traversal with 64-element stack
• Separate Plane Handling: Planes not in BVH (brute-force tested)

Future Improvements

• SAH-based splitting for better tree quality
• GPU-based construction for dynamic scenes
• Traversal optimization (persistent threads, warp-level)

References

• [Wald 2007] "On fast Construction of SAH-based BVH"
• [Aila & Karras 2010] "Understanding Ray Traversal Efficiency on GPUs"