2

在 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 个。

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

4

1 回答 1

8

首先，我想指出的是，size当您需要获取内存中布局的 Swift 类型的实际长度时，您不应该使用。你应该使用stride它。根据 Swift 的类型布局：

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

如果您想更好地理解该主题，此答案将详细介绍 Swift 中的内存布局。

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

结构的大小（在 Swift 的情况下为步幅）需要是任何结构成员的最大对齐方式的倍数。您的最大对齐Metal struct是8 个字节（对齐float2），因此在值之后 struct 的尾部有一个填充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 个字节。

同时，您的4 个字节Swift struct的最大对齐方式只需要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 中使用float2fromsimd而不是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 bytes ) 而不是size( 20 bytes ) 来获取长度。

3x3 矩阵情况也是如此：Metal 的float3x3大小为48 字节，对齐方式为16 字节。正如我假设的那样，您已经创建了一个Swift structwith 9 Floats ，它的步幅/大小为36 bytes和对齐4 bytes。因此，错位。使用matrix_float3x3从simd.

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

于 2018-02-11T00:29:58.590 回答