Get to know Metal function pointers

Highlight

Metal 函数指针是 2020 年引入的重要特性，它允许 GPU 代码在运行时通过函数表（Function Table）动态调用不同的函数，而不是在编译时硬编码调用关系。

核心内容

光线追踪里的着色阶段很容易膨胀。一个简单 path tracer 先生成射线，做几何求交，再根据命中的材质和光源计算颜色。只要场景里出现多种光源、多种材质，原本放在 shade 里的逻辑就会变成一串固定调用，新增一种材质就要重新组织 shader 代码和 pipeline。

Metal 2020 的函数指针 API 把这件事拆开了。开发者可以把光照函数、材质函数标记为 visible function，再通过 Metal API 把这些函数接入 pipeline。GPU 端拿到的是 visible function table（可见函数表），shader 可以按索引取出函数指针并调用它。

这场 session 的主线不是只给 shader 增加一个语法特性。它把函数指针放进完整的工程流程里：哪些函数参与单次 pipeline 编译，哪些函数预编译成 binary function，哪些函数后续增量加入，怎样把函数表传给 GPU，递归和 SIMD divergence 会带来哪些性能约束。

详细内容

03:09 从固定的着色调用开始

session 先用一个简化的 path tracer 展示问题位置。kernel 在求交之后找到材质，再调用 shade 完成光照和材质计算。

float3 shade(...);

[[kernel]] void rtKernel(...
                         device Material *materials,
                         constant Light &light)
{
    // ...

    device Material &material = materials[intersection.geometry_id];
    float3 result = shade(ray, triangleIntersectionData, material, light);

    // ...
}

关键点：

intersection.geometry_id 用来找到当前命中的材质。
shade 是光照和材质逻辑的汇合点。
当光源类型和材质类型增加，固定的 shade 调用会承担越来越多分支。

shade 内部再分成 lighting 和 material 两个阶段。

float3 shade(...)
{
    Lighting lighting = LightingFromLight(light, triangleIntersectionData);

    return CalculateLightingForMaterial(material, lighting, triangleIntersectionData);
}

关键点：

LightingFromLight 代表根据光源类型计算 lighting 的阶段。
CalculateLightingForMaterial 代表根据材质类型应用 lighting 的阶段。
这两个阶段正好适合拆成可替换的 visible function。

04:59 用 `[[visible]]` 暴露 GPU 函数

Metal Shading Language 增加了 [[visible]] 函数限定。它和 vertex、fragment、kernel 一样标在函数定义前，表示这个函数可以被 Metal API 操作。

[[visible]]
Lighting Area(Light light, TriangleIntersectionData triangleIntersectionData)
{
    Lighting result;

    // Clever math code ...

    return result;
}

关键点：

[[visible]] 让 Area 这样的光照函数成为可被 CPU 端引用的 Metal function object。
transcript 明确提到 visible function 可以在另一个 Metal 文件或另一个 Metal library 中。
pipeline 创建时可以把这些 function object 加进去，之后 GPU 端通过函数指针调用。

05:30 单次编译：把 visible function 复制进 pipeline

single compilation 的做法是把所有可能调用的 visible function 放入 MTLLinkedFunctions.functions，再创建 compute pipeline state。

// Single compilation configuration
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [area, spot, sphere, hair, glass, skin]
computeDescriptor.linkedFunctions = linkedFunctions

// Pipeline creation
let pipeline = try device.makeComputePipelineState(descriptor: computeDescriptor,
                                                   options: [],
                                                   reflection: nil)

关键点：

linkedFunctions.functions 列出 pipeline 可能调用的 visible function。
pipeline 创建结果包含 kernel 代码和这些函数的 specialized copies。
transcript 把这种优化类比为 link-time optimization：运行时性能最好，pipeline 创建时间和对象体积会增加。

07:43 分离编译：把函数预编译成 binary function

separate compilation 让函数以独立 GPU binary 的形式存在，多个 pipeline 可以共享这些 binary function。入口是 MTLFunctionDescriptor。

// Create by function descriptor:
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Area"
// More configuration goes here
let areaBinaryFunction = try library.makeFunction(descriptor: functionDescriptor)

关键点：

MTLFunctionDescriptor 可以给函数创建过程附加配置。
设置 name 后，library.makeFunction(descriptor:) 返回对应的 Metal function object。
descriptor 的价值在下一步：请求 binary pre-compilation。

// Create and compile by function descriptor:
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Area"
functionDescriptor.options = MTLFunctionOptions.compileToBinary
let areaBinaryFunction = try library.makeFunction(descriptor: functionDescriptor)

关键点：

MTLFunctionOptions.compileToBinary 要求 Metal 把函数预编译为 binary。
binary function 可以在 pipeline 创建时被引用，减少把函数复制进 pipeline 的成本。
session 明确说明这种模式会带来函数调用的运行时开销。

预编译函数通过 linkedFunctions.binaryFunctions 接入 pipeline。

// Specify binary functions on compute pipeline descriptor
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [spot, sphere, hair, glass, skin]
linkedFunctions.binaryFunctions = [areaBinaryFunction]
computeDescriptor.linkedFunctions = linkedFunctions

// Pipeline creation
let pipeline = try device.makeComputePipelineState(descriptor: computeDescriptor,
                                                   options: [],
                                                   reflection: nil)

关键点：

functions 里的函数会参与 specialization。
binaryFunctions 里的函数使用预编译 binary。
两者可以混用，用 pipeline size、创建时间和 runtime performance 做取舍。

11:04 增量编译：给已有 pipeline 加函数

游戏和动态内容加载常会遇到新材质。session 用 Wood material 举例，展示如何让已有 pipeline 支持后续增加 binary function。

// Create initial pipeline with option
computeDescriptor.supportAddingBinaryFunctions = true

// Create and compile by function descriptor
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Wood"
functionDescriptor.options = MTLFunctionOptions.compileToBinary
let wood = try library.makeFunction(descriptor: functionDescriptor)

// Create new pipeline from existing pipeline
let newPipeline =
   try pipeline.makeComputePipelineStateWithAdditionalBinaryFunctions(functions: [wood])

关键点：

初始 pipeline 创建前要设置 supportAddingBinaryFunctions = true。
新函数仍然通过 MTLFunctionDescriptor 和 compileToBinary 生成。
新 pipeline 从已有 pipeline 派生，只附加新增的 binary function。

12:22 visible function table：把函数指针传给 GPU

visible function table 是 GPU 端拿到函数指针的方式。shader 里先声明函数指针类型，再把 function table 作为 kernel 参数或 argument buffer 传入。

// Helper using declaration in Metal
using LightingFunction = Lighting(Light, TriangleIntersectionData);
using MaterialFunction = float3(Material, Lighting, TriangleIntersectionData);

// Specify tables as kernel parameters
visible_function_table<LightingFunction> lightingFunctions [[buffer(1)]],
visible_function_table<MaterialFunction> materialFunctions [[buffer(2)]],

// Access via index
LightingFunction *lightingFunction = lightingFunctions[light.index];
Lighting lighting = lightingFunction(light, triangleIntersection);
return materialFunctions[material.index](material, lighting, triangleIntersection);

关键点：

visible_function_table<LightingFunction> 和 visible_function_table<MaterialFunction> 分别承载光照函数和材质函数。
shader 通过 light.index、material.index 选择函数。
取出的函数指针可以先保存到变量，也可以直接从 table 调用。

CPU 端从 pipeline state 分配 table，再把 function handle 放进指定 index。

// Initialize descriptor
let vftDescriptor = MTLVisibleFunctionTableDescriptor()
vftDescriptor.functionCount = 3
let lightingFunctionTable = pipeline.makeVisibleFunctionTable(descriptor: vftDescriptor)!

// Find and set functions by handle
let functionHandle = pipeline.functionHandle(function: spot)!
lightingFunctionTable.setFunction(functionHandle, index:0)

// Find and set functions by handle
computeCommandEncoder.setVisibleFunctionTable(lightingFunctionTable, bufferIndex:1)
argumentEncoder.setVisibleFunctionTable(lightingFunctionTable, index:1)

关键点：

MTLVisibleFunctionTableDescriptor.functionCount 决定 table 容量。
pipeline.functionHandle(function:) 取得可写入 table 的 handle。
table 可以绑定到 computeCommandEncoder，也可以放进 argument buffer；使用 argument encoder 时 transcript 提醒要调用 useResource。

14:23 function groups：告诉编译器每个调用点可能调用谁

function groups 用来给编译器更多信息。shader 在调用点标记 group，CPU 端在 MTLLinkedFunctions.groups 中声明每个 group 的候选函数。

// Add function groups to our shading function
float3 shade(...)
{
    LightingFunction *lightingFunction = lightingFunctions[light.index];
    [[function_groups("lighting")]] Lighting lighting = lightingFunction(light,
triangleIntersection);

    MaterialFunction *materialFunction = materialFunctions[material.index];
    [[function_groups("material")]] float3 result = materialFunction(material, lighting, triangleIntersection);
    return result;
}

关键点：

[[function_groups("lighting")]] 标记 lighting 调用点。
[[function_groups("material")]] 标记 material 调用点。
编译器知道每个调用点的函数集合后，可以在生成 pipeline 时做更有针对性的优化。

// Function Group configuration
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [area, spot, sphere, hair, glass, skin]
linkedFunctions.groups = ["lighting" : [area, spot, sphere ],
                          "material" : [hair, glass, skin ] ]
computeDescriptor.linkedFunctions = linkedFunctions

关键点：

groups 字典把 group name 映射到该调用点可能调用的函数。
lighting group 只包含 area、spot、sphere。
material group 只包含 hair、glass、skin。

15:25 递归和 SIMD divergence 的性能边界

函数指针也打开了递归实现的空间。session 以 recursive path tracer 为例：材质函数可以发出新的 ray，再求交、着色、继续递归。为了表达调用链深度，compute pipeline descriptor 提供 maxCallStackDepth。默认值是 1，适合常见 visible function 和 intersection function 场景；如果函数调用链更深，需要显式设置预期深度。

最后一个边界是 divergence。SIMD group 最适合所有线程执行同一条指令。通过函数指针调用时，同一个 SIMD group 内的线程可能指向不同函数，最坏情况下要把这些函数依次执行。session 给出的处理思路是：把函数参数、线程 index、函数 index 写入 threadgroup memory，按函数 index 排序，再按排序结果调用函数，把结果写回 threadgroup memory，让每个线程读取自己的结果。对于复杂函数，这能降低 divergence 的额外成本。

核心启发

1. 可热更新的材质函数库

做什么：把材质 BRDF 拆成 visible function，运行时用 function table 按材质索引选择。

为什么值得做：session 明确展示了 lighting 和 material 两个可拆阶段，材质种类增加时不必把所有逻辑塞进一个 shade。

怎么开始：先把一个材质函数标记为 [[visible]]，用 MTLLinkedFunctions.functions 接入 pipeline，再用 visible_function_table<MaterialFunction> 从 GPU 端按 material.index 调用。

2. 面向流式资源的 shader 增量加载

做什么：游戏加载新资产时，把新材质函数编译为 binary function，并从已有 pipeline 派生新 pipeline。

为什么值得做：transcript 提到动态环境里常会出现新函数，尤其是 games streaming 新资产和 shader 的场景。

怎么开始：创建初始 pipeline 前设置 computeDescriptor.supportAddingBinaryFunctions = true，新材质出现后用 MTLFunctionDescriptor、compileToBinary 和 makeComputePipelineStateWithAdditionalBinaryFunctions(functions:) 补上。

3. 可配置的 lighting/material 预览器

做什么：做一个开发工具视图，让美术或图形工程师切换不同 light 和 material 组合，实时观察同一场景的变化。

为什么值得做：visible function table 可以把 lighting function 和 material function 分开传给 GPU，shader 通过索引组合调用。

怎么开始：分别建立 lightingFunctions 和 materialFunctions 两张表，CPU 端用 setFunction(_:index:) 写入候选函数，UI 改变时只更新索引或表项。

4. 递归 path tracer 的调用栈实验

做什么：在一个小型 path tracer 中比较迭代 bounce 和递归 bounce 的实现成本。

为什么值得做：session 明确指出 visible function 调用链可能需要更深 stack，ray tracing 也可以改写成迭代方式降低 stack usage。

怎么开始：先用默认 maxCallStackDepth 跑单层 visible function 调用，再逐步增加递归深度，记录 pipeline 配置和帧耗时变化。

5. 函数指针 divergence 诊断视图

做什么：记录每个 threadgroup 内的 function index 分布，识别同一 SIMD group 中函数调用高度分散的材质或光源组合。

为什么值得做：session 明确指出函数指针可能造成 thread-level divergence，复杂函数下可以通过排序增加 coherence。

怎么开始：在实验 kernel 中把参数、thread index 和 function index 写入 threadgroup memory，按 function index 排序后调用函数，再比较排序前后的耗时。

关联 Session

Discover ray tracing with Metal — 本场用 ray tracing 的 shading 阶段解释函数指针，先看这场能补齐加速结构、求交和 compute kernel 的上下文。
Build GPU binaries with Metal — 深入 shader 编译、GPU binary、binary archive 和动态库，和本场的 binary function 取舍相邻。
Debug GPU-side errors in Metal — function table、argument buffer 和 GPU 端资源绑定出错时，需要 Xcode 12 的 GPU-side error 报告能力。
Optimize Metal apps and games with GPU counters — 本场最后讨论 runtime overhead 和 divergence，这场提供用 GPU counters 定位瓶颈的工具路线。
Harness Apple GPUs with Metal — 了解 Apple GPU 和 Metal workload 的性能背景，有助于判断函数指针方案的成本边界。

Highlight

核心内容

详细内容

03:09 从固定的着色调用开始

04:59 用 [[visible]] 暴露 GPU 函数

05:30 单次编译：把 visible function 复制进 pipeline

07:43 分离编译：把函数预编译成 binary function

11:04 增量编译：给已有 pipeline 加函数

12:22 visible function table：把函数指针传给 GPU

14:23 function groups：告诉编译器每个调用点可能调用谁

15:25 递归和 SIMD divergence 的性能边界

核心启发

1. 可热更新的材质函数库

2. 面向流式资源的 shader 增量加载

3. 可配置的 lighting/material 预览器

4. 递归 path tracer 的调用栈实验

5. 函数指针 divergence 诊断视图

关联 Session

评论

04:59 用 `[[visible]]` 暴露 GPU 函数