HIP Ray Tracing 2.0 引入了改进的自定义函数处理、Bitcode 链接等功能

首次发布：2023年3月28日

最后更新：2024年6月26日

Daniel Meister

Paritosh Kulkarni

我们很高兴地宣布 HIPRT-v2 的发布，其中包含多项新功能和优化。为了克服一些效率低下的问题，我们对 API 进行了一些更改。在这篇博文中，我们将概述这些更改，并通过几个简单的示例展示新功能。

自定义函数表

我们做出的最大改变在于处理自定义函数及其在自定义函数表中的排列方式。在以前的版本中，我们使用函数指针为自定义图元传递相交函数。事实证明，函数指针在 HIP 中效率非常低下，造成了主要的瓶颈。为了绕过此限制，我们在编译之前将函数名（而非函数指针）传递给 HIPRT。自定义函数表是一个有界的结构，会在内部构建，允许编译器全面了解情况，从而完全优化代码。因此，此解决方案几乎没有开销，并且使设置更加简单。

自定义函数表是一个二维结构，其中每个条目由光线类型索引（行）和几何类型索引（列）标识。该表被展平为一个一维数组，使得第一个条目对应第一个光线类型和第一个几何类型，第二个条目对应第一个光线类型和第二个几何类型，依此类推。

让我们通过一个简单的示例来展示它是如何工作的。首先，我们需要在构造 hiprtGeometry 之前设置 hiprtGeometryBuildInput 中的 geomType，以便稍后在自定义函数表中引用它。

hiprtGeometryBuildInput geomInput;
geomInput.geomType = 0;

我们希望为自定义图元使用以下相交函数。请注意，函数签名与先前版本不同，它将命中属性封装在 hiprtHit 结构中，从而使其更紧凑。

__device__ bool intersectFunc(const hiprtRay& ray, const void* data, void* payload, hiprtHit& hit) {...}

我们在 hiprtFuncNameSet 结构中设置函数名，该结构代表自定义函数表中的一个条目。

hiprtFuncNameSet funcNameSet;
funcNameSet.intersectFuncName = "intersectFunc";

我们将条目分配到表中所需的位置。注意二维表是如何展平为一维数组的。

constexpr int RayTypes = 2;
constexpr int GeomTypes = 3;
hiprtFuncNameSet funcNameSets[RayTypes * GeomTypes];

int geomTypeIndex = 1;
int rayTypeIndex = 2;
funcNameSets[rayTypeIndex * GeomTypes + geomTypeIndex] = funcNameSet;

我们将函数名表和维度传递给跟踪内核编译函数。或者，我们也可以使用 hiprtBuildTraceKernels 而不是 hiprtBuildTraceKernelsFromBitcode。

hiprtApiFunction function;
hiprtBuildTraceKernelsFromBitcode(..., GeomTypes, RayTypes, funcNameSets, &function);

现在，我们必须创建自定义函数表本身。

hiprtFuncTable funcTable;
hiprtCreateFuncTable(hiprtContext, GeomTypes, RayTypes, &funcTable);

通常，我们希望将与自定义图元对应的数据传递给相交函数。

hiprtFuncDataSet funcDataSet;
functDataSet.intersectFuncData = ...;
hiprtSetFuncTable(hiprtContext, funcTable, geomTypeIndex, rayTypeIndex, funcDataSet);

最后，我们将表传递给遍历对象。请注意，在 HIPRT-v2 中，遍历对象的参数顺序已更改。

hiprtSceneTraversalClosest tr(scene, ray, hiprtFullRayMask, hiprtTraversalHintDefault, nullptr, table);

在 HIPRT-v2 中，除了自定义相交函数，我们还可以设置筛选函数，用于过滤一些不希望的相交。例如，自相交或 alpha 遮蔽。函数签名与相交函数非常相似，只是 hit 是常量引用，因为相交已经找到。我们希望决定是否过滤掉找到的相交（返回 true），或者将其报告为正确的命中（返回 false）。

__device__ bool filterFunc(const hiprtRay& ray, const void* data, void* payload, const hiprtHit& hit) {...}

筛选函数的名称以与相交函数名称相同的方式传递。

hiprtFuncNameSet funcNameSet;
funcNameSet.filterFuncName = "filterFunc";

编译和 Bitcode 链接

我们所做的第二个重大更改与跟踪内核编译过程有关。在早期版本中，跟踪内核代码是即时组装的，并由 HIP 运行时编译 API 进行编译。然而，HIP 中 Bitcode 链接的引入开辟了以前无法使用的设备代码链接的新可能性。我们利用这项技术将预编译的 HIPRT 遍历代码链接到用户提供的应用程序代码，从而实现了更简洁、更快的编译过程。现在，API 提供了 hiprtBuildTraceKernelsFromBitcode 函数来实现这种新方法，该函数在内部在线完成所有工作。

也可以手动进行离线 Bitcode 链接，这在某些情况下可能很有用，例如，当在线链接不可行时（由于各种限制），或者当我们希望节省链接本身所需的时间时。

我们需要以下三个文件：

用户内核代码：user_code.cpp
hiprt 编译后的 bitcode：hiprt02000_amd_lib_win.bc（HIPRT SDK 的一部分）
自定义函数表：custom_function_table.cpp（见下文）

对于以下命令，我们假设 HIP SDK 已安装在系统上，包括 HIP SDK 中的 hipcc 和 clang，并且它们已添加到 **PATH** 中。请在 Windows 上使用 Windows Shell 或 PowerShell（WSL、MinGW、cygwin 等与 **hipcc** 和 **clang** 不兼容）。

我们使用以下命令将 **user_code.cpp** 编译为 bitcode，假设我们为 Navi21（gfx1030）构建。

hipcc -O3 -std=c++17 --offload-arch=gfx1030 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm -I../../ -ffast-math -D BLOCK_SIZE=64 -D SHARED_STACK_SIZE=16 user_code.cpp -parallel-jobs=15 -o user_code.bc

请注意，仅当使用 hiprtGlobalStack 时，才需要 BLOCK_SIZE 和 SHARED_STACK_SIZE。

我们使用以下命令将用户 bitcode 与 HIPRT 库 bitcode 链接。

clang -fgpu-rdc --hip-link --cuda-device-only --offload-arch=gfx1030 user_code.bc  hiprt02000_amd_lib_win.bc -o offline_linked_user_code.hipfb

如果我们的代码中有筛选函数或相交函数，我们必须手动编写自定义函数表（**custom_function_table.cpp**）。假设我们有一个名为 myFilter 的筛选函数，自定义函数如下所示：

#if defined(__CUDACC__)
#include <cuda_runtime.h>
#include <cmath>
#else
#include <hip/hip_runtime.h>
#endif
#include <hiprt/hiprt_device.h>

__device__ bool myFilter(const hiprtRay& ray, const void* data, void* payload, const hiprtHit& hit);

__device__ bool intersectFunc(
  unsigned int geomType, unsigned int rayType, const hiprtFuncTableHeader& tableHeader, const hiprtRay& ray, void* payload, hiprtHit& hit)
{
  const unsigned int  index = tableHeader.numGeomTypes * rayType + geomType;
  const void* data  = tableHeader.funcDataSets[index].intersectFuncData;
  switch (index)
  {
  default: {
    return false;
  }
  }
}

__device__ bool filterFunc(
    unsigned int geomType, unsigned int  rayType, const hiprtFuncTableHeader& tableHeader, const hiprtRay& ray, void* payload, const hiprtHit& hit)
{
  const unsigned int  index = tableHeader.numGeomTypes * rayType + geomType;
  const void* data  = tableHeader.funcDataSets[index].filterFuncData;
  switch (index)
  {
  case 0: {
    return myfilter( ray, data, payload, hit );
  }
  default: {
    return false;
  }
  }
}

请注意光线类型和几何类型如何映射到索引（请参阅上一节）。我们可以以相同的方式添加更多自定义函数。

我们使用以下命令编译 **custom_function_table.cpp**。

hipcc -O3 -std=c++17 --offload-arch=gfx1030 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm -I../../ -ffast-math custom_function_table.cpp -parallel-jobs=15 -o custom_function_table.bc

我们使用以下命令将所有内容链接在一起。

clang -fgpu-rdc --hip-link --cuda-device-only --offload-arch=gfx1030 user_code.bc  custom_function_table.bc hiprt02000_amd_lib_win.bc -o offline_linked_user_code.hipfb

一旦我们有了链接的 HIP fat binary 文件，就可以将其加载为 HIP 模块，并从中查询必要的函数指针。

其他更改

我们为 Navi3x 引入了光线遍历提示，通过允许用户添加有关输入光线分布的额外知识（例如阴影或反射光线），从而提供对性能优化的新一级控制。可以通过遍历对象构造函数指定这些遍历提示。

hiprtSceneTraversalClosest tr(scene, ray, hiprtTraversalHintShadowRays);

为了方便用户，我们增加了对变换矩阵的支持。在以前的版本中，我们仅支持由组件（SRT）定义的 SRT 变换。现在，用户可以在 hiprtSceneBuildInput 结构中指定变换类型。

hiprtSceneBuildInput sceneInput;
sceneInput.frameType  = hiprtFrameTypeMatrix;

为了区分变换结构，我们有两种结构：用于分量表示的 hiprtFrameSRT 和用于矩阵表示的 hiprtFrameMatrix。请注意，无论类型如何，HIPRT 都能在运动模糊的情况下正确处理帧的插值。

我们可以在 hiprtRay 结构中指定最小 t 值，这对于次表面散射等情况可能很有用。请注意，时间参数已移至遍历对象构造函数的最后一个参数。

hiprtRay ray;
ray.minT = 0.1f;
hiprtSceneTraversalClosest tr(scene, ray, hiprtFullRayMask, hiprtTraversalHintDefault, nullptr, nullptr, 0, time);

我们将构造 API 设计为线程安全的，允许使用 HIP 流在多个线程中并发构造多个 hiprtGeometry 结构。
我们扩展了 API，允许用户通过传递跟踪函数数量、跟踪函数名称数组和输出函数数组来编译多个模板化的跟踪内核。

int numFunctions = 2;
const char* funcNames = {"Trace<true>", "Trace<false>"};
hiprtApiFunction functionsOut[2];
hiprtBuildTraceKernelsFromBitcode(..., numFunctions, funcNames, ..., functionsOut);