---
title: Ray Tracing with HIP on AMD GPUs
url: https://www.emergentmind.com/topics/ray-tracing-using-hip
type: topic
---

# Ray Tracing with HIP on AMD GPUs

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 [2603.00292].

## 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: $t = ( -b\pm\sqrt\Delta )/(2a)$.

- **Triangle**: For triangle vertices $(\mathbf{v}_0,\mathbf{v}_1,\mathbf{v}_2)$, using the normal $\mathbf{n}=(\mathbf{v}_1-\mathbf{v}_0)\times(\mathbf{v}_2-\mathbf{v}_0)$, the intersection parameter is
  \[
  t=\frac{\mathbf{n}\cdot(\mathbf{v}_0-\mathbf{o})}{\mathbf{n}\cdot\mathbf{d}}
  \]
  and point-in-triangle is verified with three edge tests:
  \[
  \begin{aligned}
    \mathbf{n}\cdot[(\mathbf{v}_1-\mathbf{v}_0)\times(\mathbf{p}-\mathbf{v}_0)]&\geq0 \\
    \mathbf{n}\cdot[(\mathbf{v}_2-\mathbf{v}_1)\times(\mathbf{p}-\mathbf{v}_1)]&\geq0 \\
    \mathbf{n}\cdot[(\mathbf{v}_0-\mathbf{v}_2)\times(\mathbf{p}-\mathbf{v}_2)]&\geq0
  \end{aligned}
  \]
  with $\mathbf{p} = \mathbf{o}+t\mathbf{d}$. 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 $\mathtt{min}=(x_{\min},y_{\min},z_{\min})$ and $\mathtt{max}=(x_{\max},y_{\max},z_{\max})$. Ray–AABB intersection is performed in $O(1)$ via the slab method.
- **BVH Construction**: Typically binary or $4$–$8$-ary; leaf nodes store primitives, internal nodes encapsulate child AABBs.
- **Split Strategy**: The Surface Area Heuristic (SAH) guides node splitting to minimize
  \[
  C = C_{\mathrm{trav}} + \frac{\mathrm{SA}(L)}{\mathrm{SA}(N)}C(L) + \frac{\mathrm{SA}(R)}{\mathrm{SA}(N)}C(R)
  \]
  where $\mathrm{SA}(\cdot)$ is surface area. The standard top-down recursive build (see pseudocode in the data) partitions primitives at each node by choosing a split minimizing $C$.
- **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**:
  ```cpp
  hiprtContextCreationInput input{};
  input.deviceType=hiprtDeviceAMD;
  input.ctxt = oroGetRawCtx(oroCtx);
  input.device=oroGetRawDevice(device);
  hiprtContext ctx;
  hiprtCreateContext(HIPRT_API_VERSION,input,ctx);
  hiprtSetLogLevel(hiprtLogLevelWarn|hiprtLogLevelError);
  ```

- **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 $64$ threads, thus block dimensions should be multiples of $64$.
- **Memory Coalescing**: Structuring AABB and triangle data as Structure-of-Arrays, $16$-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 $2\times$ faster traversal relative to naive BVH structures.
  - On AMD MI100, throughput of roughly $300$ million primary rays per second is achieved with $1$ 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 [2603.00292].

Source: https://www.emergentmind.com/topics/ray-tracing-using-hip