Discover ray tracing with Metal

Highlight

Metal 2020 的光线追踪 API 把 ray intersector 带进 Metal Shading Language，让开发者在单个 compute kernel 中完成射线生成、求交和着色，并用 intersection functions 自定义透明材质和自定义几何体的命中逻辑。

---

核心内容

在这场 session 之前，Apple 已经讲过用 Metal Performance Shaders 做光线追踪。那条路径能工作：先启动一个 compute kernel 生成 rays，把它们写入 Metal buffer，再用 MPSRayIntersector 求交，最后启动另一个 compute kernel 读取 intersection result 并做 shading。每次光线继续弹射，应用还要把 rays 和 intersections 在 buffer 之间来回传递。（01:33）

2020 年的 Metal ray tracing API 改变了这个编程模型。intersector 现在可以直接在 Metal Shading Language 里创建和调用，原先分散在多个 kernel 和中间 buffer 里的工作，可以合并进一个 compute kernel。这样做减少了 ray 和 intersection 数据的内存读写，也让多次弹射的外层循环可以留在 shader 代码里表达。（02:14）

这套 API 的另一个核心是 acceleration structure。开发者提供 triangle 或 bounding box 这样的几何描述，Metal 负责构建用于快速求交的数据结构。新的控制点在于：应用可以决定 acceleration structure 的内存何时分配、scratch buffer 如何创建、构建命令排在哪个 GPU command queue 和 command buffer 上。（03:48）

session 的后半段处理更棘手的场景：透明材质和自定义几何体。alpha test 树叶如果每次遇到透明像素都重新发射一条 ray，会不断从 acceleration structure 根节点重新遍历，成本很高。Metal 的 triangle intersection function 可以在遍历过程中接受或拒绝一次命中；bounding box intersection function 则可以把 sphere、curve、hair strand 这类自定义 primitive 接入同一套求交流程。（10:01）

Apple 的 demo 也给了一个现实尺度：Advanced Content Team 用这套 API 做了一个 path tracing 场景，编辑时用每像素一个 sample 做交互预览，离线 fly-through 用每像素 400 个 samples 收敛到干净图像；场景来自 Quixel Megascans，超过 3200 万个 triangles。（08:28）

---

详细内容

在 compute kernel 里直接求交

（02:42）基础 kernel 的形状很直接：每个像素生成一条 ray，创建 intersector<triangle_data>，把 ray 和 primitive_acceleration_structure 传进去，拿到 intersection_result<triangle_data> 后继续 shading。

[[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...
}

关键点：

• primitive_acceleration_structure 是 kernel 的普通 buffer binding，后面用 compute command encoder 绑定。
• ray 来自当前像素的相机数学，session 没展开 generateCameraRay 的实现。
• intersector<triangle_data> 在 shader 中直接创建，省掉了 MPS 方案中的 ray buffer 和 intersection buffer 往返。
• intersection_result<triangle_data> 交给后续 shading；如果 shading 太复杂，session 建议用 profiling 决定是否拆分 kernel。

（07:25）CPU 侧绑定 acceleration structure 也走专门的 encoder API。

computeEncoder.setAccelerationStructure(accelerationStructure, bufferIndex: 0)

关键点：

• bufferIndex: 0 对应 kernel 参数上的 [[buffer(0)]]。
• acceleration structure 不是普通 MTLBuffer，但绑定模型和 buffer binding point 对齐。
• transcript 还提到 argument encoder 也有对应绑定方法。

构建 primitive acceleration structure

（04:48）构建从 descriptor 开始。这里使用 triangle geometry，把 vertex buffer 和 triangle count 填进 MTLAccelerationStructureTriangleGeometryDescriptor。

let accelerationStructureDescriptor = MTLPrimitiveAccelerationStructureDescriptor()

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

geometryDescriptor.vertexBuffer = vertexBuffer
geometryDescriptor.triangleCount = triangleCount

accelerationStructureDescriptor.geometryDescriptors = [ geometryDescriptor ]

关键点：

• MTLPrimitiveAccelerationStructureDescriptor 描述 primitive acceleration structure。
• MTLAccelerationStructureTriangleGeometryDescriptor 描述 triangle geometry。
• transcript 明确说 triangle geometry 可以有自己的 vertex buffer、index buffer、triangle count 等信息；这个 snippet 展示了最小路径。
• triangle 数据会嵌入 acceleration structure，构建完成后不必为了 acceleration structure 持有原 vertex buffer。

（05:46）Metal 先让应用查询需要的大小，再由应用分配 acceleration structure 和 scratch buffer。

// 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)!

关键点：

• device.accelerationStructureSizes 返回 acceleration structure 和 build scratch buffer 的大小。
• makeAccelerationStructure 让应用控制 allocation 的时间和位置。
• scratch buffer 只在 build 期间作为临时存储使用。
• session 选择 .storageModePrivate，因为 CPU 不需要访问这段临时数据。

（06:24）真正的 build 被编码进 GPU command buffer。

// 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()

关键点：

• makeAccelerationStructureCommandEncoder 创建新的 acceleration structure encoder。
• build 命令拿到 descriptor、目标 acceleration structure 和 scratch buffer。
• transcript 说明 build 运行在 GPU timeline 上，没有 CPU synchronization。
• build 后可以安全地在 GPU 上继续排 intersection work。

用 triangle intersection function 做 alpha test

（12:16）alpha test 的低效点在于透明命中。如果每次透明都重新发 ray，遍历会从 acceleration structure 根节点重新开始。triangle intersection function 可以在遍历过程中检查 alpha texture，并决定这次 intersection 是否接受。

[[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;
}

关键点：

• [[intersection(triangle, triangle_data)]] 声明这是 triangle intersection function，并要求访问 barycentric coordinate。
• primitiveIndex、geometryIndex 和 barycentricCoords 来自 intersector。
• function 可以绑定自己的 resource，这里用 materials 读取 alpha texture 和 UV。
• 返回 true 接受这次命中，返回 false 让 traversal 继续寻找下一个候选命中。

用 bounding box function 接入自定义 primitive

（14:38）自定义 primitive 的入口是 bounding box。应用先提供包围盒，Metal 对这些 boxes 建 acceleration structure；真正的 sphere、curve 或 hair strand 求交，由 bounding box intersection function 负责。

// 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 ]

关键点：

• descriptor 仍然是 MTLPrimitiveAccelerationStructureDescriptor。
• geometry descriptor 换成 MTLAccelerationStructureBoundingBoxGeometryDescriptor。
• boundingBoxBuffer 存放 axis-aligned bounding boxes。
• Metal 会先找 ray 和 bounding box 的潜在命中，再调用应用提供的 intersection function。

（15:38）sphere 示例展示了 bounding box intersection function 的职责：数学求交、检查距离范围、返回是否接受以及 intersection distance。

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

[[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 };
}

关键点：

• BoundingBoxResult 用 attribute 标出是否接受和命中距离。
• origin、direction、minDistance、maxDistance 是 intersector 传入的 ray 数据和范围。
• function 自己用 intersectRaySphere 做数学测试；session 没展开这个 helper 的实现。
• 返回的 distance 会参与整个场景中最近 intersection 的比较。

（16:20）如果 shading 需要额外数据，可以用 ray payload 把 intersection function 的结果带回 compute kernel。

[[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 };
}

关键点：

• ray_data float3 & normal [[payload]] 是 payload reference。
• intersection function 写入 normal 后，发起求交的 kernel 可以读到它。
• transcript 提醒 payload 的修改在接受和拒绝 intersection 时都会可见，通常只应在准备接受命中时写 payload。

连接 intersection functions 和 function table

（17:18）intersector 要调用 intersection functions，必须先把这些 functions 链接进 compute pipeline state。

// 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)

关键点：

• intersection function 从 Metal library 加载，方式和普通 Metal function 一样。
• MTLLinkedFunctions 把这些 functions 附到 compute pipeline descriptor。
• 编译 pipeline 时，compiler 会把 compute kernel 和 intersection functions 链接到一起。
• 这一步让 compute kernel 内部的 ray intersector 可以调用自定义 intersection function。

（18:35）链接后，还要用 intersection function table 把 function 映射到具体 geometry 或 instance。

// 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)

关键点：

• MTLIntersectionFunctionTableDescriptor 决定 table 容量。
• table 从具体 computePipeline 创建，只能和这个 pipeline 或派生 pipeline 一起使用。
• functionHandle 指向已链接进 pipeline state 的 executable code。
• function table 也是 resource 绑定点，示例把 sphere buffer 绑定到 index 0。

（19:26）kernel 使用 function table 时，把它作为另一个 buffer 参数传给 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...
}

关键点：

• intersection_function_table<triangle_data> 是 shader 端看到的 table 类型。
• acceleration structure 和 function table 分别位于 [[buffer(0)]]、[[buffer(1)]]。
• intersector 根据 geometry descriptor 和 instance descriptor 上的 table offset 选择要调用的 entry。
• CPU 侧用 encoder.setIntersectionFunctionTable(functionTable, bufferIndex: 1) 绑定 table。

---

核心启发

1. 交互式 path tracing 预览器

做什么：给 3D 场景编辑器加一个 noisy preview，用户移动相机时每像素发一条 ray，停下后逐步累积 samples。

为什么值得做：session 的 demo 正是这个工作流；交互编辑看速度，离线渲染看收敛质量。

怎么开始：先按 rtKernel 的结构做 one thread per pixel，把 ray generation、intersection、shading 放进同一个 compute kernel，再增加 sample accumulation。

2. 支持透明树叶和镂空材质的 ray traced shadow

做什么：在光追阴影里正确处理 alpha texture，让树叶、栅栏、网格布料不会被当成完整 triangle。

为什么值得做：triangle intersection function 可以在 acceleration structure traversal 中直接 reject 透明命中，避免每次透明都重新发 ray。

怎么开始：为带 alpha texture 的 geometry 设置 intersection function table entry，在 function 中用 barycentric coordinate 插值 UV，采样 alpha 后返回 accept/reject。

3. 用数学 primitive 渲染 sphere、curve 或 hair

做什么：把球体、曲线、头发丝这类对象作为 bounding box primitive 接入 renderer，减少三角化带来的内存和外观问题。

为什么值得做：session 展示了 sphere 和 cubic Bezier hair strand 的路径；curve 可以保持平滑，hair strand 只需要少量 control points。

怎么开始：为每个 primitive 生成 axis-aligned bounding box，使用 MTLAccelerationStructureBoundingBoxGeometryDescriptor 建 acceleration structure，再写 [[intersection(bounding_box)]] function。

4. 给 renderer 加 kernel 拆分实验开关

做什么：在同一个 renderer 中保留单 kernel 和多 kernel 两条路径，用实际场景测试 occupancy、memory traffic 和 frame time。

为什么值得做：session 明确提醒 intersection 很重；如果 shading 也很复杂，合并后的 compute kernel 可能降低 occupancy。

怎么开始：先实现单 kernel 路径，再把 heavy shading 分离成另一个 compute pass，用 GPU counters 或 Xcode Metal tools 对比。

5. 把 acceleration structure 构建放进资源调度系统

做什么：让场景加载器决定 acceleration structure 何时分配、构建、refit 或 compact，避免渲染循环里临时管理大块内存。

为什么值得做：新 API 给应用控制 allocation 和 GPU timeline build 的位置；session 还提到 instancing、refitting 和 compaction 可用于动态几何和内存回收。

怎么开始：围绕 device.accelerationStructureSizes、makeAccelerationStructure、scratch buffer 和 makeAccelerationStructureCommandEncoder 封装一个 build job。

---

关联 Session

• Get to know Metal function pointers — 本场结尾直接推荐的配套 session，继续解释 arbitrary material code、dynamic shader call 和 ray tracing function table。
• Harness Apple GPUs with Metal — 补齐 Apple GPU 与 Metal 渲染架构背景，适合理解高性能 renderer 的基础约束。
• Optimize Metal apps and games with GPU counters — 对应本场关于 profiling 和 kernel 拆分的提醒，用 GPU counters 查找 Metal workload bottleneck。
• Gain insights into your Metal app with Xcode 12 — 介绍 Xcode 12 的 Metal debugger 和性能建议，适合排查 ray tracing renderer 的资源、带宽和 API 使用问题。
• Bring your Metal app to Apple silicon Macs — 讲 Apple silicon Mac 上的 Metal/TBDR 架构和移植问题，适合把图形密集型 Metal app 带到新 Mac 平台。