Enie
Enie

Reputation: 699

structs with arrays in Metal

My goal is to create a struct with containing arrays in swift that I can pack into a metal buffer to use in a kernel function.

something like this:

struct StructA  {
  var listA: [StructB] = [StructB](repeating: StructB(), count: 100)
  var listB: [StructC] = [StructC](repeating: StructC(), count: 100)
  var primitiveA: Int = 0
  var primitiveB: Int = 0
}

I'm not quite sure how to solve this. The example above does obviously not work since the arrays are not physically inside StructA. My guess is there must be a way to create an array inside a struct so that the memory used is physically aligned inside the struct.

Currently my workaround is to pack listA and listB into two separate buffers, set these as parameters of my kernel function and within the kernel function assign them to a StructA that is created in each thread, which is a crazy redundancy.

kernel void functionA(
  const device StructB *listA [[buffer(0)]],
  const device StructC *listB [[buffer(1)]],
  device int &primitiveA [[buffer(2)]],
  device int &primitiveB [[buffer(3)]],
) {
  StructA structA = StructA(); //create a struct in each and every thread
  structA.listA = listA;       //and assign the same lists to the struct
  structA.listB = listB;       //that is the same in every thread
  structA.primitiveA = primitiveA;
  structA.primitiveB = primitiveB;

  //do stuff with structA
}

The example might not be flawless, but I think the problem is described sufficiently. I hope there is a solution to this. If not by creating sufferable structs, I would also use any other solution that solves the redundancy.

Upvotes: 3

Views: 913

Answers (1)

Palle
Palle

Reputation: 12109

Doing this would require the arrays listA and listB to be allocated on the stack, which is currently not possible, as arrays don't have a fixed size and are therefore allocated on the heap.

To solve this, you could use tuples instead of arrays or declare your struct in C code. You could even share the struct declaration with your metal code by putting them in a C header file.

This answer may help you.

Upvotes: 1

Related Questions