近两天,AMD 发布了 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 中用于硬件光线追踪的光线加速器。
特点
在第一个版本中,AMD 实现了光线追踪应用程序所需的一些基本功能。
三角形网格是 HIP RT 的基本原语,AMD 从中构建 hiprtGeometry。通常,场景由许多网格组成。开发人员可以利用这些网格之上的 hiprtScene 支持来创建一个加速结构,当光线与许多网格相交时,该结构可用于有效地搜索网格。此过程允许开发人员找到最近的光线命中或捕获所有命中。第二种情况对阴影和透明度很有用。
为了加速光线追踪操作,HIP RT 在库中构建了加速结构。没有适合所有情况的单一结构,因此,AMD 提供多种选择以找到适合您需求的结构。例如,您可以选择一个高质量的构建,它比其他构建需要更长的时间,但结果是一个高质量的数据结构,它可以使单个光线追踪查询更快。或者,高级光线追踪用户可以选择使用 HIP RT 来加载自建的包围体层次 (BVH),然后可以使用 HIP RT 在 GPU 上执行相交。时至今日,这是 HIP RT 的一项独特功能,其他光线追踪 API 都不支持。例如,HIP RT 可用于为场景中的静态几何体构建高质量的 BVH,然后在运行时加载它。
加速构建器不仅从三角形网格中获取三角形,而且还获取轴对齐边界框 (AABB)。如果在 GPU 上执行更改逻辑时对象正在动态更改,则此过程很有用。这方面的一个例子是剥皮。由于缓冲区是通过 AABB 传递的,因此数据永远不会离开 GPU。用户只需在 GPU 上执行几何更新逻辑,将更新的 AABB 写入缓冲区,然后将其传递给构建器。
如果您的场景包含三角形以外的图元,HIP RT 也可以处理它们。您需要做的就是为几何体提供一个自定义的交集函数,在 HIP RT 中注册该函数,然后当射线与图元相交时将调用此函数。
你的虚拟相机的曝光时间是有限的吗?然后,您可能想尝试我们的运动模糊。您需要做的就是在曝光期间设置一些变换。关键帧的间隔不需要统一。您可以在曝光时间内传递关键帧的位置。
代码示例
现在大家已经基本了解了 HIP RT 的功能,让我们来看看 API。
以下是一些基本示例,但教程中提供了更多详细信息。
创建 Context
创建 HIP 上下文和设备后,将它们传递hiprtCreateContext
给创建 HIP RT 上下文。您可以在 00_context_creation 找到代码。
创建 Geometry
创建hiprtContext
完成后,下一步就是构建hiprtGeometry
我们需要在设备端执行光线交集的数据结构。创建hiprtGeometry
步骤如下:
hiprtTriangleMeshPrimitive mesh; //set up the mesh //get the temporary buffer size hiprtGetGeometryBuildTemporaryBufferSize( ... ); //create hiprtGeometry hiprtGeometry geom; hiprtCreateGeometry( ... ); //build hiprtGeometry hiprtBuildGeometry( ... );
首先创建一个hiprtTriangleMeshPrimitive
对象,然后hiprtGeomety
通过调用创建一个hiprtCreateGeometry()
可以内置的对象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 hiprtScene hiprtCreateScene( ... ); //build hiprtScene hiprtBuildScene( ... );
尽管输入与 building 不同hiprtGeometry
,但结构是相同的。现在我们需要通过将场景信息存储到传递给()的hiprtSceneBuildInput
对象来传递场景信息。hiprtBuildScene
示例代码可以在 02_scene_intersection 中找到。
在设备端,我们可以重用上面的程序。我们需要做的就是获取hiprtScene
并创建一个hiprtSceneTraversalClosest
对象。然后,您应该能够在场景中追踪光线并渲染这样的图像(当然是我们最喜欢的!)。
高级示例
虽然我们在这里只展示了非常基础的内容,但您可以在我们的 SDK 中找到一些有趣的教程,例如使用共享堆栈以获得更好的性能、运动模糊和自定义交集函数。
下载 HIP RT
在这篇文章中,我们介绍了 HIP RT(HIP 中的光线追踪库),并快速浏览了 API。如果你觉得这个第一个版本对您的项目有用。您可以期待未来 HIP RT 的改进