Passing parameters in Metal Compute Kernel using Swift 4

Question

On the CPU side, I have a structure I want to pass to the compute kernel:

  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
  }

Before running the kernel, I pass in the data to the MTLComputeCommandEncoder:

Option 1 (directly):

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

Option 2 (indirectly via MTLBuffer):

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

Either option works fine if the 'dummy' variable exists in the structure, but fails if the 'dummy' variable does not exists. The code fails in the call to:

commandEncoder!.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)

With the error:

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

On the Metal Kernel Side, here's the relevant code snippets:

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);
}

I also get a similar issue when trying to pass in a 3x3 matrix to the compute kernel. It complains that 36 bytes are provided but expecting 48.

Anyone got any ideas on the issue?

Answer 1

The float2 type in the Metal structure requires an alignment of 8 bytes. The stride of the struct then becomes 24 so every value is aligned properly if you use it in an array. Your Swift structure only uses float which requires an alignment of 4 bytes, and an overall stride of 20.

This presentation by Apple goes into some detail (starting at 30:20 of the video). The Metal Shading Language Specification lists the alignment of vector and matrix data types in sections 2.2 and 2.3.

Similarly, float3x3 has a size of 48 and an alignment of 16 in Metal. But an array of 9 floats would only have a size of 36 and an alignment of 4. Using the simd module can provide you data types in Swift with the same size and alignment requirements as those used in Metal.

As a side note, you should be using MemoryLayout<>.stride instead of MemoryLayout<>.size when you call setBytes. The following Swift structure has a stride of 24 but a size of only 20.

struct BoundingBoxParameters {
    var topLeft: float2 = float2(0)   // float2 is from "import simd"
    var size: float2 = float2(0)
    var levelOfDetail: Float = 1.0
}

Answer 2

Utilize simd to synchronize your alignment:

import simd

private struct BoundingBoxParameters {

    let topLeft: float2 = float2(0.0, 0.0)
    let size: float2 = float2(0.0, 0.0)
    let levelOfDetail: Float = 1.0

    init(topLeft: float2,
         size: float2,
         levelOfDetail: Float) {
        self.topLeft = topLeft
        self.size = size 
        self.levelOfDetail = levelOfDetail
    }
}