Ray Tracing with HIP on AMD GPUs

• Ray tracing using HIP is a technique for photorealistic rendering that couples AMD's HIP programming model with HIPRT's hardware-accelerated ray tracing.
• It employs robust acceleration structures like BVH and AABB to optimize primitive intersection calculations and enhance scalability.
• Advanced memory management and hardware-specific optimizations on AMD GPUs yield significant performance improvements, achieving up to 2× faster traversal.

Ray tracing using HIP refers to the implementation of high-performance ray-tracing algorithms on AMD GPUs through the Heterogeneous-compute Interface for Portability (HIP), with additional acceleration via the HIP Ray Tracing (HIPRT) framework. HIP provides a platform-agnostic, C++-like API for direct GPU programming, while HIPRT exposes hardware raytracing support in newer AMD architectures. The workflow covers foundational geometric computations, BVH construction and traversal, and device-specific optimization techniques for scalable photorealistic rendering (Yoshimura et al., 27 Feb 2026).

1. Geometric Fundamentals of Ray Tracing

Ray tracing is founded on mathematical formalisms for modeling ray–primitive intersections. A ray is defined parametrically as

$$\mathbf{r}(t) = \mathbf{o} + t\,\mathbf{d}\,, \quad t\ge0$$

where $\mathbf{o}$ is the origin and $\mathbf{d}$ the direction (typically unit length).

Primitive Intersection:

• Sphere: For a sphere centered at $\mathbf{c}$ with radius $R$, the condition

$\|\mathbf{x}-\mathbf{c}\|^2 = R^2$

yields a quadratic in $t$ when substituting $\mathbf{x}=\mathbf{o}+t\mathbf{d}$. The discriminant

$\Delta = b^2-4ac$

with

$$a = \mathbf{d}\cdot\mathbf{d},\quad b = 2\,\mathbf{d}\cdot(\mathbf{o}-\mathbf{c}),\quad c = \|\mathbf{o}-\mathbf{c}\|^2 - R^2$$

determines the presence of real intersections: $\mathbf{o}$0.

• Triangle: For triangle vertices $\mathbf{o}$1, using the normal $\mathbf{o}$2, the intersection parameter is

$\mathbf{o}$3

and point-in-triangle is verified with three edge tests:

$\mathbf{o}$4

with $\mathbf{o}$5. Listing 1 in the source report provides the device implementation.

2. Hierarchical Acceleration Structures

Efficient ray tracing depends critically on robust acceleration structures. The principal data structure is the Bounding Volume Hierarchy (BVH).

• AABB: Each node comprises $\mathbf{o}$6 and $\mathbf{o}$7. Ray–AABB intersection is performed in $\mathbf{o}$8 via the slab method.
• BVH Construction: Typically binary or $\mathbf{o}$9–$\mathbf{d}$0-ary; leaf nodes store primitives, internal nodes encapsulate child AABBs.
• Split Strategy: The Surface Area Heuristic (SAH) guides node splitting to minimize

$\mathbf{d}$1

where $\mathbf{d}$2 is surface area. The standard top-down recursive build (see pseudocode in the data) partitions primitives at each node by choosing a split minimizing $\mathbf{d}$3.

• Traversal: Closest-hit traversal evaluates intersections recursively, always descending into children whose AABBs intersect the ray, visiting leaf nodes to test primitives directly.

3. HIP Programming Model and Memory Semantics

HIP exposes explicit control over GPU memory and kernel launch configuration.

• Memory Allocation and Transfer:
  • Allocate device memory: hipMalloc(&d_ptr, size)
  • Copy to device: hipMemcpy(d_ptr, h_ptr, size, hipMemcpyHostToDevice)
  • Free device memory: hipFree(d_ptr)
• Kernel Dispatch and Synchronization:
  • Create streams: hipStreamCreate(&stream)
  • Launch kernels: hipLaunchKernelGGL(kernel,..., stream,...)
  • Synchronize: hipStreamSynchronize(stream) or hipDeviceSynchronize()
• Device Query:
  • Get device count: hipGetDeviceCount(&n)
  • Get properties: hipGetDeviceProperties(&props,dev)

Orochi wrappers may be used for raw context/device queries, relevant when interoperating with HIPRT.

4. HIPRT Framework and API Workflows

HIPRT provides high-performance ray tracing on AMD hardware, abstracting BVH construction and leveraging the ray accelerator units on RDNA2/3 GPUs.

• Context Creation:

$\mathbf{c}$0

• BVH Building:
  • Bottom-level: Fill hiprtTriangleMeshPrimitive with device pointers to vertex/index arrays, set build input type, query buffer size, create and build geometry.
  • Top-level: Define hiprtInstance array (hiprtInstanceTypeGeometry), provide transforms, upload with oroMalloc and oroMemcpyHtoD, then create and build scene BVH.
• Custom Primitives:
  • Use hiprtAABBListPrimitive with user AABB arrays; implement custom intersection in device code and register via hiprtFuncTable/hiprtFuncDataSet.
• Traversal and Stack Management:
  • Shared and global stack buffers enable efficient traversal across threads and thread blocks (detailed in Listing 3). Support exists for both any-hit (shadow rays, early termination) and closest-hit traversals.
• Kernel Compilation and Invocation:
  • Use hiprtBuildTraceKernels with a hiprtFuncNameSet referencing custom functions, producing an hiprtApiFunction invocable as a standard HIP kernel.

5. End-to-End Pipeline and Use Cases

The stepwise workflow proceeds as follows:

• On the host, allocate geometry and parameter buffers.
• Instantiate the HIPRT context, build geometry/scene BVHs.
• Compile the ray tracing kernel with hiprtBuildTraceKernels.
• Launch the kernel per pixel or sample, generating primary rays, traversing BVHs for intersections, and shading as appropriate for direct lighting or a full path-tracing loop.
• Copy result buffers back to host for image output.

The architecture supports rendering triangles, spheres, and arbitrary custom primitives. The use of the any-hit device traversal accelerates shadow ray computation, while the closest-hit routine enables path tracing, next-event estimation, and related algorithms.

6. AMD Hardware Optimizations

Several best practices directly target the architectural features of AMD GPUs:

• Workgroup Sizing: For RDNA2/3, optimal performance is observed at wavefronts of $\mathbf{d}$4 threads, thus block dimensions should be multiples of $\mathbf{d}$5.
• Memory Coalescing: Structuring AABB and triangle data as Structure-of-Arrays, $\mathbf{d}$6-byte aligned, ensures maximal memory throughput.
• Hardware Acceleration: HIPRT auto-targets the DPS units; setting buildFlags=PreferBalancedBuild results in near-optimal runtime vs. build-time trade-off.
• Benchmarks:
  • HIPRT balanced BVHs yield approximately $\mathbf{d}$7 faster traversal relative to naive BVH structures.
  • On AMD MI100, throughput of roughly $\mathbf{d}$8 million primary rays per second is achieved with $\mathbf{d}$9 million triangles.

7. Best Practices and Troubleshooting

Robust operational guidelines include:

• Monitoring HIP/HIPRT error codes using hipGetErrorString.
• Setting log levels (hiprtLogLevelWarn|Error) for diagnostic insight during BVH build and traversal.
• Retaining temporary buffers until all build operations complete.
• For dynamic scenes, updating the top-level BVH while marking bottom-level BVHs as static enhances update efficiency.
• Employing any-hit traversal for shadow computations, which terminate at the first occluding primitive.
• For deeper studies: “HIPRT: A Ray Tracing Framework in HIP”, “A Survey on Bounding Volume Hierarchies for Ray Tracing”, and “Physically Based Rendering” are foundational references.

This comprehensive HIPRT pipeline enables researchers to implement HIP-only ray tracers spanning triangle meshes, explicit primitives, and hybrid schemes, directly utilizable for modern physically based rendering and interactive applications (Yoshimura et al., 27 Feb 2026).