使用 Swift 4 在 Metal Compute Kernel 中传递参数

Question

在 CPU 方面，我有一个想要传递给计算内核的结构：

  private struct BoundingBoxParameters {
    var x: Float = 0
    var y: Float = 0
    var width: Float = 0
    var height: Float = 0
    var levelOfDetail: Float = 1.0
    var dummy: Float = 1.0  // Needed for success
  }

在运行内核之前，我将数据传递给MTLComputeCommandEncoder：

选项1（直接）：

commandEncoder!.setBytes(&params, length: MemoryLayout<BoundingBoxParameters>.size, index: 0)

选项 2（间接通过 MTLBuffer）：

boundingBoxBuffer.contents().copyBytes(from: &params, count: MemoryLayout<BoundingBoxParameters>.size)
commandEncoder!.setBuffer(boundingBoxBuffer, offset: 0, index: 0)

如果结构中存在“dummy”变量，则任何一个选项都可以正常工作，但如果“dummy”变量不存在则失败。代码调用失败：

commandEncoder!.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)

出现错误：

validateComputeFunctionArguments:820: failed assertion `Compute Function(resizeImage): argument params[0] from buffer(0) with offset(0) and length(20) has space for 20 bytes, but argument has a length(24).'

在 Metal Kernel 方面，这里是相关的代码 sn-ps：

struct BoundingBoxParameters {
  float2 topLeft;
  float2 size;
  float levelOfDetail;
};

kernel void resizeImage(constant BoundingBoxParameters *params [[buffer(0)]],
                        texture2d<half, access::sample> sourceTexture [[texture(0)]],
                        texture2d<half, access::write> destTexture [[texture(1)]],
                        sampler samp [[sampler(0)]],
                        uint2 gridPosition [[thread_position_in_grid]]) {
  float2 destSize = float2(destTexture.get_width(0), destTexture.get_height(0));
  float2 sourceCoords = float2(gridPosition) / destSize;
  sourceCoords *= params->size;
  sourceCoords += params->topLeft;
  float lod = params->levelOfDetail;
  half4 color = sourceTexture.sample(samp, sourceCoords, level(lod));
  destTexture.write(color, gridPosition);
}

在尝试将 3x3 矩阵传递给另一个计算内核时，我也遇到了类似的问题。它抱怨提供了 36 个字节，但预期为 48 个。

有人对这个问题有任何想法吗？

Answer 1

首先，我想指出的是，当您需要获取内存中布局的 Swift 类型的实际长度时，您不应该使用 size。您应该为此使用stride。根据斯威夫特的Type Layout：

最终的大小和对齐方式是聚合的大小和对齐方式。类型的步幅是最终大小四舍五入到对齐。

This answer 详细介绍了 Swift 中的内存布局，如果您想更好地理解该主题。

---

问题是您的 Metal struct 使用 float2 和 Swift struct 用两个单独的 Float 字段替换它具有不同的内存布局。

结构的大小（在 Swift 的情况下为步幅）需要是任何结构成员的最大对齐的倍数。 Metal struct 中的最大对齐是 8 字节（float2 的对齐），因此在 float 值之后的结构尾部有一个填充。

struct BoundingBoxParameters {
    float2 topLeft; // 8 bytes
    float2 size; // 8 bytes
    float levelOfDetail; // 4 bytes
    // 4 bytes of padding so that size of struct is multiple 
    // of the largest alignment (which is 8 bytes)

}; // 24 bytes in total

所以你的 Metal struct 实际上最终会占用 24 个字节，正如错误所暗示的那样。

同时，你的Swift struct，最大对齐4字节，只需要20字节。

private struct BoundingBoxParameters {
    var x: Float = 0 // 4 bytes
    var y: Float = 0 // 4 bytes
    var width: Float = 0 // 4 bytes
    var height: Float = 0 // 4 bytes
    var levelOfDetail: Float = 1.0 // 4 bytes
    // no need for any padding 

} // 20 bytes in total

这就是为什么它们最终彼此不兼容，dummy 字段将 4 个缺失字节 补偿到 Swift struct。

为了解决这个问题，我建议你在 Swift 中使用 simd 中的 float2 而不是 Floats：

import simd 

private struct BoundingBoxParameters {
    var topLeft = float2(x: 0, y: 0)
    var size = float2(x: 0, y: 0)
    var levelOfDetail: Float = 1.0 
}

不要忘记使用MemoryLayout<BoundingBoxParameters>.stride（24 字节）而不是size（20 字节）来获取长度.

---

对于 3x3 矩阵情况也是如此：Metal 的 float3x3 的大小为 48 字节，对齐方式为 16 字节。正如我所假设的，您已经创建了一个带有 9 个Floats 的Swift struct，其步幅/大小为36 字节，对齐方式为4 字节。因此，错位。使用来自simd 的matrix_float3x3。

一般来说，当你在 Metal 中使用向量或矩阵时，你应该在 Swift 中使用对应的 simd 类型。