2
$\begingroup$

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?

$\endgroup$
2
$\begingroup$

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
}
$\endgroup$

Your Answer

By clicking "Post Your Answer", you acknowledge that you have read our updated terms of service, privacy policy and cookie policy, and that your continued use of the website is subject to these policies.

Not the answer you're looking for? Browse other questions tagged or ask your own question.