使用 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).'

在金属内核方面，这里是相关的代码片段：

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

首先，我想指出你不应该使用 size when you need to get an actual length of a Swift type laid out in memory. You should use stride for that. According to Swift's Type Layout:

The final size and alignment are the size and alignment of the aggregate. The stride of the type is the final size rounded up to alignment.

如果您想更好地理解该主题，

会在 Swift 中详细介绍内存布局。

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

结构的大小（在 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 个 Float 的 Swift struct，其中 stride/size 的 36 字节 和 [=62= 的对齐方式]4 字节。因此，错位。使用 simd 中的 matrix_float3x3。

一般来说，对于在 Metal 中使用向量或矩阵的任何情况，都应该在 Swift 中使用相应的 simd 类型。

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

Passing parameters in Metal Compute Kernel using Swift 4

kernel

image-processing

metal

metalkit

swift4