今天,我们发布了 HIP RT – 一个新的 HIP 光线追踪库。HIP RT 使在 HIP 中编写光线追踪应用程序变得容易,该库和 API 的设计宗旨是易于使用并能轻松集成到任何现有的 HIP 应用程序中。
虽然有一些其他光线追踪 API 引入了许多新方面,但我们通过消除学习许多新内核类型的需求,以略有不同的方式设计了 HIP RT。HIP RT 引入了 hiprtGeometry 和 hiprtScene 等新对象类型。一旦将几何信息传递给 HIP RT,该库就会构建数据结构,然后将其传递给 HIP 内核。在此阶段,可以使用设备端库 API 来执行相交测试。
当前一代显卡,例如基于 AMD RDNA™ 2 架构的 GPU,支持硬件光线追踪加速以进一步优化渲染时间。然而,到目前为止,支持 HIP 的应用程序还无法利用此硬件加速。HIP RT 的设计目的是允许开发人员充分利用 AMD RDNA 2 架构 GPU 中用于硬件光线追踪的光线加速器。
功能
在此第一个版本中,我们实现了一些光线追踪应用程序所需的が基本功能。
三角形网格是 HIP RT 的基本图元,我们从中构建 hiprtGeometry。通常,场景由许多网格组成。开发人员可以利用这些网格之上的 hiprtScene 支持来创建加速结构,当光线与许多网格相交时,可以使用该结构来高效地搜索网格。此过程允许开发人员找到光线的最近命中点或捕获所有命中点。后一种情况对于阴影和透明度很有用。
为了加速光线追踪操作,HIP RT 在库中构建了一个加速结构。没有一个结构适合所有情况,因此,我们提供了几种选项供您选择,以找到适合您需求的结构。例如,您可以选择一个高质量的构建,它需要比其他选项稍长的时间来生成,但结果是一个高质量的数据结构,可以加快单次光线追踪查询的速度。或者,高级光线追踪用户可以选择使用 HIP RT 来加载自建的包围盒层次结构(BVH),然后 HIP RT 可以用于在 GPU 上执行相交。如今,这是 HIP RT 的一个独特功能,其他光线追踪 API 均不支持。例如,HIP RT 可用于为场景中的静态几何体构建高质量的 BVH,然后在运行时加载它。
我们的加速构建器不仅接受来自三角形网格的三角形,还接受轴对齐包围盒(AABBs)。如果对象在动态变化,此过程很有用,因为更改逻辑在 GPU 上执行。一个例子是蒙皮。将一个包含 AABBs 的缓冲区传递时,数据永远不会离开 GPU。用户只需在 GPU 上执行几何更新逻辑,将更新后的 AABBs 写入缓冲区,然后将其传递给我们即可。
如果您的场景包含三角形以外的其他图元,HIP RT 也可以处理它们。您需要做的就是为图元提供自定义相交函数,在 HIP RT 中注册该函数,然后当光线与图元相交时将调用该函数。
您的虚拟相机的曝光时间是有限的吗?那么,您可能想尝试我们的运动模糊。您需要做的就是在曝光时间内设置一些变换。关键帧之间的时间间隔不需要是均匀的。您可以在曝光时间内传递关键帧的位置。
代码示例
既然您对 HIP RT 的功能有了初步了解,让我们来看看 API。
这里有一些基本示例,但更详细的信息可在教程中找到。
创建上下文
创建 HIP 上下文和设备后,将其传递给 hiprtCreateContext 以创建 HIP RT 上下文。您可以在 00_context_creation 中找到代码。
构建几何体
创建 hiprtContext 后,下一步是构建 hiprtGeometry,这是我们在设备端执行光线相交所需的数据结构。创建 hiprtGeometry 的步骤如下:
hiprtTriangleMeshPrimitive mesh;//set up the mesh//get the temporary buffer sizehiprtGetGeometryBuildTemporaryBufferSize( ... );//create hiprtGeometryhiprtGeometry geom;hiprtCreateGeometry( ... );//build hiprtGeometryhiprtBuildGeometry( ... );您首先创建一个 hiprtTriangleMeshPrimitive 对象,然后通过调用 hiprtCreateGeometry() 创建一个 hiprtGeomety,可以在 hiprtBuildGeometry() 中构建它。您可以在 01_geom_intersection 中找到示例代码。
上面的代码用于启动设备端的光线投射过程,所以现在让我们转向设备代码。我们的第一个内核可以非常简单:
#include <hiprt/hiprt_device.h>
extern "C"__global__ void MeshIntersectionKernel(hiprtGeometry geom, unsigned char* gDst, int2 cRes){ const int gIdx = blockIdx.x * blockDim.x + threadIdx.x; const int gIdy = blockIdx.y * blockDim.y + threadIdx.y;
hiprtRay ray; ray.origin = { gIdx / (float)cRes.x, gIdy / (float)cRes.y, -1.0f}; ray.direction = { 0.0f, 0.0f, 1.0f}; ray.maxT = 1000.f;
hiprtGeomTraversalClosest tr(geom, ray); hiprtHit hit = tr.getNextHit(); bool hasHit = tr.hasHit();
int dstIdx = gIdx + gIdy * cRes.x; gDst[ dstIdx * 4 + 0 ] = hasHit ? ((float)gIdx / cRes.x) * 255 : 0; gDst[ dstIdx * 4 + 1 ] = hasHit ? ((float)gIdy / cRes.y) * 255 : 0; gDst[ dstIdx * 4 + 2 ] = 0; gDst[ dstIdx * 4 + 3 ] = 255;}显然,要使用我们的内置函数,您需要包含 hiprt_device.h。然后您可以看到我们正在接收一个 hiprtGeometry 对象,该对象刚刚在主机端构建。要执行光线场景相交,请填充 hiprtRay,它只是光线原点和方向,然后创建 hiprtGeomTraversalClosest,并调用 getNextHit(),这将返回命中信息。就这么简单!
构建场景
如果场景中只有一个三角形网格,构建 hiprtGeometry 就足够了。然而,在大多数情况下,场景中有许多对象,我们也会考虑这些对象。因此,我们有 hiprtScene 对象,它将这些 hiprtGeometry 对象捆绑在一起,然后创建一个我们可以传递给设备端的单个对象。hiprtScene 的创建如下。
hiprtScene scene;//create hiprtScenehiprtCreateScene( ... );//build hiprtScenehiprtBuildScene( ... );虽然输入与构建 hiprtGeometry 不同,但结构是相同的。现在我们需要通过将其存储到 hiprtSceneBuildInput 对象中来传递场景信息,该对象被传递给 hiprtBuildScene()。您可以在 02_scene_intersection 中找到示例代码。
在设备端,我们可以重用上面的程序。我们所需要做的就是获取 hiprtScene 并创建一个 hiprtSceneTraversalClosest 对象。然后您应该能够通过场景追踪光线并渲染出这样的图像(当然是我们最喜欢的一个!)。

高级示例
虽然我们这里只展示了最基本的内容,但在我们的 SDK 中可以找到一些有趣的教程,例如用于提高性能的共享堆栈使用、运动模糊和自定义相交函数。


立即下载 HIP RT!
在这篇博文中,我们介绍了 HIP RT,这是一个 HIP 中的光线追踪库,并快速浏览了 API。我们希望您觉得这个第一个版本对您的项目有用。您可以在未来期待 HIP RT 的改进,敬请关注。
访问我们新的 HIP RT 页面立即下载