image-processing - 使用 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)

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

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 方面，这里是相关的代码片段:

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 当您需要获取内存中布局的 Swift 类型的实际长度时。你应该使用 stride 为了那个原因。根据 Swift 的 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.

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

---

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

结构的大小(在 Swift 的情况下为步长)需要是任何结构成员的最大对齐的倍数。 Metal struct 中最大的对齐方式是 8 个字节(float2 的对齐方式)所以在 float 之后的结构尾部有一个填充值(value)。

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 .

要解决此问题，我建议您使用 float2来自 simd在 Swift 中而不是 Float小号:

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 字节。正如我假设的那样，您已经创建了一个 Swift struct与 9 Float s 的步幅/大小为 36 字节，对齐方式为 4 字节。因此，错位。使用 matrix_float3x3来自 simd .

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