利用 Metal 探索光线追踪

利用 Metal 探索光线追踪

使用光线追踪技术，在你的 app 与游戏中展现照片一般逼真的 3D 场景——这是 Metal 图形框架与 Shading Language 的核心部分之一。探索 Metal 光线追踪 API 与 Shading Language 光线追踪扩展的基础功能，了解如何在你的图形 app 与游戏中使用它们，并掌握如何控制多个内核，将它们组合成一整个计算内核，获取最佳性能。

资源
相关视频

WWDC23
- Metal 光线追踪指南
WWDC21
- 探索 Metal 调试、性能分析及资产创建工具
下载

（你好 WWDC 2020）
你好欢迎来到 WWDC （利用 Metal 探索光线追踪）嗨我叫 Sean James 我是 GPU 软件工程师在本次演讲中我们会谈到今年我们为 Metal 引入的全新光线追踪 API 这个新的 API 让大家可以直接从任何内核函数中执行求交测试它还能让大家编写求交函数以便自定义求交过程光线追踪的应用是基于追踪光线与场景交互时经过路径光线追踪在很多领域都有应用比如声音模拟和物理模拟但是最常见的应用是真实感渲染在渲染中光线追踪用于模拟每一道光线这样我们就能模拟反射、柔化阴影间接光照等效果我们先回顾光线追踪的原理然后我们再看新的 Metal 光线追踪 API
我们先谈生成光线光线从相机发射到场景中然后我们测试光线与场景中几何形的相交每一个交点代表着光经过一个表面后反弹有多少光线反弹并且反弹方向如何能决定物体的外观我们会计算每个交点的颜色并更新图像这个过程称为着色它也可以生成额外的光线进一步反弹到场景中我们对这些光线也进行求交测试并且重复这一过程重复的次数随我们的意愿而定以便模拟光线在场景中反弹在过去两年里我们已经讲过怎么通过 Metal Performance Shaders 框架来做这个工作我们首先启动一个内核函数生成光线的初始集然后把它们写入 Metal 缓冲区然后我们可以利用 MPSRayIntersector 类使光线和场景相交求交器将求交结果写入另一个缓冲区然后我们就可以启动最终的内核函数这会读取求交数据并进行着色额外的光线可以写回光线缓冲区然后我们循环启动这两个内核函数以模拟光线反弹这是一个好方法不过它要求我们将代码拆分为三个单独的内核函数还需要我们经由内存传递光线和求交数据相反在新的 Metal 光线追踪 API 中这个求交器对象可以直接在着色语言中使用
因此我们可以将三个内核函数组成一个内核函数由于我们不需要将光线数据在内核函数之间传输我们还省去了光线及求交数据缓冲区
因此不需要靠读写内存来传递光线和求交数据这也是更为灵活的编程模型例如我们不需要反复启动内核函数外层循环可以用着色语言中的简单循环来表示我们来看看基本的光线追踪内核是怎样运行的我们启动一个二维内核函数每个线程负责一个像素我们首先生成一道光线这就是一些为给定像素点生成初始光线的计算然后我们创建求交器对象这个对象负责在光线和场景中的几何形之间寻找相交处我们可以给这个对象设置各种属性以自定义它的行为例如对于阴影和环境光遮蔽我们可以通过接受第一相交处来加速而不是穷举式地搜索最近的相交处接着我们将光线传递到求交器然后得到求交结果最后我们可以利用求交结果来进行着色请注意求交步骤要进行大量工作所以如果你们将它与复杂的着色代码结合那么你们的内核函数运行效率可能会偏低你们需要对 app 做性能分析找到最佳的办法来拆分工作分给每一个内核处理在基础求交测试中我们要做的就是这些求交函数的另一个参数就是加速结构加速结构是用于加速光线追踪过程的数据结构这些数据结构递归式地划分空间这样我们就能快速地消除不可能和给定光线相交的三角形和 MPS 一样 Metal 也能为大家代劳创建这种数据结构大家要做的仅仅是提供几何形它的运作方式是这样的首先我们需要创建一个描述符对象它描述我们要构建的加速结构然后我们需要分配内存以储存这个加速结构最后我们来构建该加速结构
我们先谈谈描述符 Metal 支持两种加速结构即图元加速结构和实例加速结构在这个示例中我们要创建一个图元加速结构顾名思义它包含了图元比如三角形图元加速结构由独立的几何形组成
Metal 还支持两类几何形现在我们重点讲三角形每个三角几何形都有自己的顶点缓冲、索引缓冲还有三角形计数等等我们首先创建一个图元加速结构描述符接着我们创建一个几何描述符在这个示例中我们只创建单一的三角形几何描述符然后我们附加顶点缓冲并定义三角形计数这个三角形数据会嵌入到加速结构中因此我们一旦建好加速结构后就不必保留顶点缓冲了然后我们将几何描述符加到加速结构描述符的一个数组中现在我们创建好了描述符我们可以为加速结构分配内存了这可能会分配掉大量内存在新的 Metal 光线追踪 API 中我们做的其中一个更新是 Metal 可以让大家全权控制加速结构分配的时间和位置这意味着大家可以重复使用内存或在 app 需要内存时分配我们首先调用 Metal 设备上的一个方法它会返回一个结构体包含我们需要分配给加速结构的内存大小接着我们就真正从 Metal 设备给加速结构分配内存最后我们还需要给暂存缓冲区分配内存 Metal 会在构建加速结构时把这个缓冲区作为临时存储空间构建完毕后就可以除去这个缓冲区我们不需要访问这个缓冲区的数据因此我们要使用私有资源储存模式以获得最佳性能我们为加速结构分配好了内存但这时加速结构还没有构建起来最后一步就是构建加速结构
在这里我们做了一些改进构建加速结构需要一些时间所以 Metal 可以让大家全权决定什么时候进行构建包括决定在哪个 GPU 指令队列和哪个指令缓冲区构建首先我们创建一个指令缓冲然后在其中创建新的 Acceleration- StructureCommandEncoders 和其它 Metal 编码器一样该对象用于调度 GPU 的工作在这个示例中我们来编码实际的建构指令提供描述符、加速结构和暂存缓冲区
最后我们结束编码提交指令缓冲这样 GPU 就可以开始构建加速结构加速结构的构建现在完全在 GPU 的时间线运行不需要 CPU 同步所以在这个指令缓冲区提交后在 GPU 中调度求交任务是安全的我们已经看过如何使用加速结构我们只要把数据传到求交器即可大家可以看到加速结构会绑定至通常的 Metal 绑定点上有一个对应的方法位于计算指令编码器和参数编码器中可用来绑定加速结构到这些绑定点上以上就是新的 Metal 光线追踪 API 的基础内容还有很多高阶的话题我们今天没有时间一一涵盖所以我推荐大家查看我们的说明文档和代码样例我还推荐大家回顾之前关于光线追踪的两个演讲因为很多概念是一样的比如 Metal 加速结构和 MPS 一样支持基元加速结构的双层次实例化这可以用于减少内存使用量 Metal 还支持整改现有的加速结构可将其与实例化结合以支持动态几何体最后 Metal 加速结构还支持压缩当加速结构建成以后可以用来回收大量内存下面我们要谈如何利用求交函数自定义求交过程但首先让我们看一个演示它涉及了我们讲到的所有内容
这是我们 Advanced Content 团队构建的一个应用情况它展示的是新的 Metal 光线追踪 API 应用于复杂的场景中我们看到的是这个场景的互动式预览它是用配备新款 AMD Radeon Pro W5700X GPU 的 Mac Pro 录制的这个互动式预览有很多噪点因为每个像素只有1个采样点但当相机停止移动时它开始聚合
这是艺术家进行编辑时看到的样子对于高质量的离线渲染我们会给每一帧充足的时间让它们聚合成无噪点的图像为了让大家看看最终的效果我们渲染了一段场景的空中俯视动画每个像素有 400 个采样点
在这里特别感谢 Quixel 公司我们使用了他们的 Megascans 资源库这个场景是由 Megascans 资源组合而成的有超过三千两百万个三角形
这个应用的运作机制和我之前提过的大纲非常相似它使用一种叫路径追踪的渲染算法来模拟高级效果比如柔化阴影、景深和间接光照
我们在互动式应用和完整的离线渲染中看过了新的 Metal 光线追踪 API 实战下面我们来谈谈新 Metal API 最让人兴奋的新功能之一那就是求交函数目前我们的编程模型是光线进入求交器然后输出相交数据但求交器却如黑箱一样看不到其中奥妙但在少数情况下大家可能想要自定义求交器的工作形式（无透明度测试有透明度测试）其中一个例子就是透明度测试这是光栅化中的常用技法利用一张材质去遮罩三角形上应该为透明的部分这可以用来增加大量几何细节而不会增加实际三角型数量这个例子里如果想让叶子得到更丰富的细节只要打开透明度测试即可
在光栅化时应用这种技法很简单只要从碎片着色器上丢弃该碎片即可但是在使用光线追踪时要高效地应用这种技法就要稍微复杂一点每次我们投射光线求交器会返回最近的相交处但是在这个情况下交点实际上是透明的因此我们要忽略它
因此现在我们要投射一条新的光线从第一个交点处发起最终光线撞上不透明的表面然后我们就可以停下来进行着色步骤了但每一条这样的光线都需要重新输入一次求交器这样做效率很低因为我们在不断地重新从树状结构的根节点开始遍历加速结构有一种更高效的应用这种技法的方法就是使用三角形求交函数
要弄清楚其原理我们需要先看一看求交器内部的工作原理
首先我们要遍历加速结构直到我们找到一处光线和一个几何形的相交处然后我们检查该相交处是否比上一个距离最近的相交处更近如果是则将其作为新的最近相交处然后我们回去继续遍历加速结构
最终我们完成这个过程并返回到最近的相交处求交函数则会直接插入这个过程每找到一个相交处这些函数都会被调用之后求交函数可以选择要接受还是拒绝该相交处若接受它将如往常一样成为新的最近相交处但它也可以拒绝该相交处这时我们会返回去继续遍历加速结构并丢弃这个相交处那么对于我们的透明度测试样例我们可以把透明度测试的逻辑移动到求交函数之中然后我们就只需要遍历一次加速结构这样效率就高了许多基本的求交函数看起来像这样首先我们声明函数 triangle 代表着这是一个三角形求交函数 triangle_data 代表我们要访问相交点的重心坐标该函数返回一个布尔值代表着要接受还是拒绝该相交我们从求交器获得相关信息比如图元索引、几何形索引还有相交点的重心坐标及其他此处用不到的数据我们还可以绑定自己的资源并用这些资源来访问透明度材质首先我们先查找这个图元的透明度材质下一步我们使用重心坐标以插入材质的坐标这和大家做着色步骤时的做法完全相同然后我们进行实际的材质查询最后我们根据交点是否是透明的来返回真或假但是求交函数不仅仅能应用于三角形几何形 Metal 同样允许你写出边界框的求交函数这些函数调用的时机是在一条光线与你提供的一个边界框相交之时这可以用来为自定图元建模在这个场景中这个球形上就应用了边界框的求交函数光线会通过数学计算来测试与球面的相交而不是使用三角形建模这是一个非常简单的例子但是边界框求交函数还可以用来给曲线这样的更复杂的表面进行建模这个例子里我们使用了边界框求交函数来渲染头发每一根头发都是自定的三次贝赛尔曲线图元我们使用了一万个这种图元来为头发塑造总体的外形（三角形贝塞尔曲线）虽然为每一个曲线图元求交会带来更大的开销但使用曲线替代三角形使得我们可以渲染出完美柔滑的绺状头发而没有任何三角形失真我们使用的内存也更少因为每一条曲线都仅仅需要几个控制点而非大量的三角形对于头发数量大的角色内存的高效使用很重要我们假设有一个场景是用球形图元生成的
我们首先要提供与轴对齐的边界框来封闭每一个球然后就像处理三角形一样 Metal 为框建立加速结构然后当我们提供光线时 Metal 寻找边界框与光线的相交处与处理三角形时一样 Metal 会在每一次发现可能的相交处时调用边界框求交函数我们只需要做一点小的改动即可构建这种加速结构
我们要创造图元加速结构描述符就像之前那样但是这一次我们不是创造三角几何形描述符而是边界框几何形描述符
这个描述符仅仅是提供了一个附加缓冲的位置其包括我们的边界框以及边界框的数量计数
就像三角形求交函数一样边界框求交函数返回值为真或假其代表接受或者拒绝该相交处但它们同样也负责计算从光线原点到自定义图元上的任意相交点的最短距离 Metal 会使用该距离来计算总体上的最近相交处范围遍及场景之中所有不同的三角形和边界框图元该声明稍有变化以显示出我们在写的是边界框求交函数由于我们需要返回多个值因此我们返回一个用户定义的结构体这个例子里我们创建的结构体包括一个布尔值用来指示要接受还是拒绝该相交处还包括一个浮点数包括相交处的距离和三角形求交函数相似我们会收到光线数据如光线原点和方向以及和相交处有关的信息比如图元索引这里我们又一次可以绑定自己的资源用它来传递一个包含球形的数组我们首先用数学方法检查光线是否会与边界框内封闭的球相交如果不相交就直接返回假下一步与三角形求交函数不同我们还有一项额外责任那就是检查求交函数是否位于可接受的值域内如果是就可以接受这个相交处并返回相交处距离但如果我们想要返回附加数据该怎么办呢？举个例子我们可能想返回相交点处的表面法线以用于着色实现这一点我们可以使用光线荷载它是一条很小的信息可以从一个求交函数处一直传递到启动了求交过程的内核函数处我们只要声明一个对我们荷载类型的引用将其声明为求交函数的参数然后我们就可以对其写入数据来改变该荷载注意对于荷载的改动是可见的这无关于你是否会接受该相交处所以大部分情况下你应该只希望在接受该相交处时才修改该荷载要取得荷载的值只要修改我们对求交器的调用来通过引用传递荷载任何由求交函数对荷载的改动都会在求交器返回值时可见目前为止我们已经讲了如何从内核函数使用光线求交器我们也讲到了如何写求交函数下面我们来讲讲如何将两者结合起来为了让求交器能够调用我们的求交函数我们需要在一条计算管线状态中将它们链接起来
我们首先把求交函数从 Metal 库中载入这和载入其他 Metal 函数相同
下面我们把求交函数接入到管线描述符通过使用 MTLLinkedFunctions 对象
最后我们如平时一样编译管线状态编译器会将内核函数的代码与求交函数链接到一起使得求交函数可以在内核函数内部由光线求交器调用现在我们把求交函数链接到了我们的计算管线状态内部我们需要将这些求交函数映射到独立的几何形中去
这样的话我们可以使用求交函数表这个表中包括了指向管线状态中实际可执行代码的指针
就如同每一条加速结构几何形可以拥有自己的顶点缓冲一样它们各自也可以拥有自己的求交函数如果大家在使用一个实例加速结构每个实例也都可以拥有自己的求交函数集
加速结构几何形与实例描述符两者都具有独立的求交函数表偏移当光线撞击一个几何形时这两个值就会相加以便决定调用求交函数表的哪一项要创建一个求交函数表我们首先创建一个描述符这个示例里我们创建一个够大的描述符足够容下我们所有的求交函数下一步我们来分配该求交函数表这个表指向特定计算管线状态中的实际可执行代码所以它也是从该管线状态中创建而来更进一步讲它只能用于该管线状态中或者用于由它衍生的管线状态中下一步我们为每一个求交函数从计算管线状态中获得一个句柄这个句柄代表了已经与管线状态链接的函数可执行代码的地址
然后我们把这个句柄插入到求交函数表内我们也是在这里绑定表中的求交函数使用的资源
比如说我们把球的缓冲绑定到 0 号缓冲索引使用一个求交函数表和使用加速结构是很相似的
我们依然只要把它传递给求交器求交函数表也会绑定到通常的 Metal 缓冲绑定点上而且也依然有一个对应的 API 方法来把一个求交函数表绑定到计算命令编码器和参数编码器上要建立求交函数就只需要我们做这么多求交函数是强大的工具可以让大家对相交过程拥有更多的掌控我总结一下我们首先讲到了全新的 Metal 光线追踪 API 以及光线求交器对象它现在已经可以在任何一个内核函数内直接使用这是一个更为灵活的编程模型让大家可以将工作以最适合自己应用的方式分割开
我们还讲到了新的 API 是如何让大家能更好地掌控加速结构在何时何处进行分配和构建
最后我们谈到了可以怎样去写求交函数以便自定义求交过程关于求交函数我们能谈的还有很多实际上我们还让这些潜在的机制也可以在其他应用中使用
比如说因为光线可以在任意时间撞击任意表面着色步骤就需要能够按需求为不时输入的临时物料执行代码所以我想建议大家观看一下函数指针方面的演讲那个演讲和这个演讲可以搭配起来大家可以从中了解另一个强大的工具它也能做同样的事情（了解 Metal 函数指针）我们今天的时间只够讲一讲新光线追踪 API 的基础所以我还推荐大家去看一看本次演讲附带的样例代码和说明文档谢谢观看祝大家享受 WWDC 之后的内容

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure [[buffer(0)]],
              /* ... */)
{
    // Generate ray
    ray r = generateCameraRay(tid);

    // Create an intersector
    intersector<triangle_data> intersector;

    // Intersect with scene
    intersection_result<triangle_data> intersection;

    intersection = intersector.intersect(r, accelerationStructure);

    // shading...
}

4:48 - Create an acceleration structure descriptor

let accelerationStructureDescriptor = MTLPrimitiveAccelerationStructureDescriptor()

// Create geometry descriptor(s)
let geometryDescriptor = MTLAccelerationStructureTriangleGeometryDescriptor()

geometryDescriptor.vertexBuffer = vertexBuffer
geometryDescriptor.triangleCount = triangleCount

accelerationStructureDescriptor.geometryDescriptors = [ geometryDescriptor ]

5:46 - Allocate acceleration storage

// Query for acceleration structure sizes
let sizes = device.accelerationStructureSizes(descriptor: accelerationStructureDescriptor)

// Allocate acceleration structure
let accelerationStructure =
    device.makeAccelerationStructure(size: sizes.accelerationStructureSize)!

// Allocate scratch buffer
let scratchBuffer = device.makeBuffer(length: sizes.buildScratchBufferSize,
                                      options: .storageModePrivate)!

6:24 - Build acceleration structure

// Create command buffer/encoder
let commandBuffer = commandQueue.makeCommandBuffer()!
let commandEncoder = commandBuffer.makeAccelerationStructureCommandEncoder()!

// Encode acceleration structure build
commandEncoder.build(accelerationStructure: accelerationStructure,
                     descriptor: accelerationStructureDescriptor,
                     scratchBuffer: scratchBuffer,
                     scratchBufferOffset: 0)

// Commit command buffer
commandEncoder.endEncoding()
commandBuffer.commit()

7:15 - Pass acceleration structure to ray intersector

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure [[buffer(0)]],
              /* ... */)
{
    // generate ray, create intersector...

    intersection = intersector.intersect(r, accelerationStructure);

    // shading...
}

7:25 - Bind acceleration structure with compute command encoder

computeEncoder.setAccelerationStructure(accelerationStructure, bufferIndex: 0)

12:16 - Triangle intersection functions

[[intersection(triangle, triangle_data)]]
bool alphaTestIntersectionFunction(uint primitiveIndex        [[primitive_id]],
                                   uint geometryIndex         [[geometry_id]],
                                   float2 barycentricCoords   [[barycentric_coord]],
                                   device Material *materials [[buffer(0)]])
{
    texture2d<float> alphaTexture = materials[geometryIndex].alphaTexture;

    float2 UV = interpolateUVs(materials[geometryIndex].UVs,
        primitiveIndex, barycentricCoords);

    float alpha = alphaTexture.sample(sampler, UV).x;

    return alpha > 0.5f;
}

14:38 - Creating a bounding box acceleration structure

// Create a primitive acceleration structure descriptor
let accelerationStructureDescriptor = MTLPrimitiveAccelerationStructureDescriptor()

// Create one or more bounding box geometry descriptors:
let geometryDescriptor = MTLAccelerationStructureBoundingBoxGeometryDescriptor()

geometryDescriptor.boundingBoxBuffer = boundingBoxBuffer
geometryDescriptor.boundingBoxCount = boundingBoxCount

accelerationStructureDescriptor.geometryDescriptors = [ geometryDescriptor ]

15:29 - Bounding Box Result

struct BoundingBoxResult {
    bool accept [[accept_intersection]];
    float distance [[distance]];
};

15:38 - Bounding box intersection functions

[[intersection(bounding_box)]]
BoundingBoxResult sphereIntersectionFunction(float3 origin            [[origin]],
                                             float3 direction         [[direction]],
                                             float minDistance        [[min_distance]],
                                             float maxDistance        [[max_distance]],
                                             uint primitiveIndex      [[primitive_id]],
                                             device Sphere *spheres   [[buffer(0)]])
{
    float distance;

    if (!intersectRaySphere(origin, direction, spheres[primitiveIndex], &distance))
        return { false, 0.0f };

    if (distance < minDistance || distance > maxDistance)
        return { false, 0.0f };

    return { true, distance };
}

16:20 - Ray payload

[[intersection(bounding_box)]]
BoundingBoxResult sphereIntersectionFunction(/* ... */,
                                             ray_data float3 & normal [[payload]])
{
    // ...

    if (distance < minDistance || distance > maxDistance)
        return { false, 0.0f };

    float3 intersectionPoint = origin + direction * distance;
    normal = normalize(intersectionPoint - spheres[primitiveIndex].origin);

    return { true, distance };
}

16:48 - Ray payload 2

[[kernel]]
void rtKernel(/* ... */)
{
    // generate ray, create intersector...

  float3 normal;

    intersection = intersector.intersect(r, accelerationStructure, functionTable, normal);

    // shading...
}

17:18 - Linking intersection functions

// Load functions from Metal library
let sphereIntersectionFunction = library.makeFunction(name: “sphereIntersectionFunction”)!
// other functions...

// Attach functions to ray tracing compute pipeline descriptor
let linkedFunctions = MTLLinkedFunctions()

linkedFunctions.functions = [ sphereIntersectionFunction, alphaTestFunction, ... ]

computePipelineDescriptor.linkedFunctions = linkedFunctions

// Compile and link ray tracing compute pipeline
let computePipeline = try device.makeComputePipeline(descriptor: computePipelineDescriptor,
                                                     options: [],
                                                     reflection: nil)

18:17 - Intersection function table offsets

class MTLAccelerationStructureGeometryDescriptor : NSObject {

    var intersectionFunctionTableOffset: Int

// ...

}

struct MTLAccelerationStructureInstanceDescriptor {
    var intersectionFunctionTableOffset: UInt32
    // ...
};

18:35 - Creating an intersection function table

// Allocate intersection function table
let descriptor = MTLIntersectionFunctionTableDescriptor()

descriptor.functionCount = intersectionFunctions.count

let functionTable = computePipeline.makeIntersectionFunctionTable(descriptor: descriptor)

for i in 0 ..< intersectionFunctions.count {
    // Get a handle to the linked intersection function in the pipeline state
    let functionHandle = computePipeline.functionHandle(function: intersectionFunctions[i])

    // Insert the function handle into the table
    functionTable.setFunction(functionHandle, index: i)
}

// Bind intersection function resources
functionTable.setBuffer(sphereBuffer, offset: 0, index: 0)

19:26 - Pass intersection function table to ray intersector

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure   [[buffer(0)]],
              intersection_function_table<triangle_data> functionTable [[buffer(1)]],
              /* ... */)
{
    // generate ray, create intersector...

    intersection = intersector.intersect(r, accelerationStructure, functionTable);

    // shading...
}

19:33 - Bind intersection function table with compute command encoder
```
encoder.setIntersectionFunctionTable(functionTable, bufferIndex: 1)
```

正在查找特定内容？在上方输入一个主题，就能直接跳转到相应的精彩内容。

提交你查询的内容时出现错误。请检查互联网连接，然后再试一次。

资源

相关视频

WWDC23

WWDC21