MSL: thread group local variables require default constructors for Metal matrices

[billhollings]

SPIR-V conversion to MSL fails compilation if a struct containing a Metal matrix is declared as a threadgroup local variable, because it requires a default constructor on all its members, and Metal matrices do not seem to have a default constructor.

program_source:129:22: error: call to implicitly-deleted default constructor of 'threadgroup _152'
    threadgroup _152 _154;
                     ^
program_source:11:14: note: default constructor of '_152' is implicitly deleted because field '_m1' has no default constructor
    float3x2 _m1;
             ^

These CTS test fail as a result:

dEQP-VK.memory_model.shared.basic_types.3
dEQP-VK.memory_model.shared.basic_types.9
dEQP-VK.memory_model.shared.basic_arrays.4
dEQP-VK.memory_model.shared.basic_arrays.6
dEQP-VK.memory_model.shared.basic_arrays.7
dEQP-VK.memory_model.shared.basic_arrays.9
dEQP-VK.memory_model.shared.arrays_of_arrays.0
dEQP-VK.memory_model.shared.arrays_of_arrays.2
dEQP-VK.memory_model.shared.arrays_of_arrays.4
dEQP-VK.memory_model.shared.arrays_of_arrays.5
dEQP-VK.memory_model.shared.arrays_of_arrays.8
dEQP-VK.memory_model.shared.nested_structs.0
dEQP-VK.memory_model.shared.nested_structs.1
dEQP-VK.memory_model.shared.nested_structs.2
dEQP-VK.memory_model.shared.nested_structs.3
dEQP-VK.memory_model.shared.nested_structs.5
dEQP-VK.memory_model.shared.nested_structs.6
dEQP-VK.memory_model.shared.nested_structs.8
dEQP-VK.memory_model.shared.nested_structs.9
dEQP-VK.memory_model.shared.nested_structs_arrays.0
dEQP-VK.memory_model.shared.nested_structs_arrays.1
dEQP-VK.memory_model.shared.nested_structs_arrays.2
dEQP-VK.memory_model.shared.nested_structs_arrays.3
dEQP-VK.memory_model.shared.nested_structs_arrays.4
dEQP-VK.memory_model.shared.nested_structs_arrays.6
dEQP-VK.memory_model.shared.nested_structs_arrays.7
dEQP-VK.memory_model.shared.nested_structs_arrays.9

SPIR-V file: matrix-default-constructor.spv.zip

[cdavis5e]

Not only that, but Metal methods have to have an address space on them. It's probably because the matrix types lack a constructor specifically for the threadgroup AS.

Wait, aren't these tests for VK_KHR_vulkan_memory_model, which we don't even support? I'm not sure we even can support it, which will likely become a problem in the future because there's talk of requiring it for later versions of Vulkan.

[billhollings]

Digging into this now.

aren't these tests for VK_KHR_vulkan_memory_model

Hmmm...I don't think so. The "memory_model" in the test names seems to be just referencing testing OpMemoryModel. And the SPIR-V here includes:

OpMemoryModel Logical GLSL450

which, AFAIK, would be set to VulkanKHR instead of GLSL450 if VK_KHR_vulkan_memory_model/SPV_KHR_vulkan_memory_model was enabled.

[billhollings]

It's probably because the matrix types lack a constructor specifically for the threadgroup AS.

This definitely appears to be the issue. I've tried numerous attempts at initializing the matrix, culminating in...

threadgroup _152 _154 = { float4(1.0, -5.0, -9.0, -5.0), float3x2(float2(1.0, -7.0), float2(1.0, 2.0), float2(8.0, 7.0)), bool4(false, true, false, false) };

which is initializing it to the values the shader is trying to set it to later. But all such attempts to initialize the matrix results in variations of:

program_source:129:62: error: no matching constructor for initialization of 'threadgroup metal::float3x2' (aka 'threadgroup matrix<float, 3, 2>')
    threadgroup _152 _154 = { float4(1.0, -5.0, -9.0, -5.0), float3x2(float2(1.0, -7.0), float2(1.0, 2.0), float2(8.0, 7.0)), bool4(false, true, false, false) };

which, unless I am missing something simple, leads me to expect that matrixes are simply not allowed in the threadgroup AS, although I can find no documentation to that effect.

Unless I AM missing something simple...options include either replacing threadgroup matrices with our own custom type (like with safe arrays?), or adding a portability flag to indicate and test for workgroupMatrices support and direct CTS tests accordingly.

[cdavis5e]

which, unless I am missing something simple, leads me to expect that matrixes are simply not allowed in the threadgroup AS, although I can find no documentation to that effect.

I think it's an oversight. There are methods that accept parameters of matrix type in the threadgroup AS. If this weren't allowed, there would be no reason for such methods. For this reason, I've filed FB9719349.

Unless I AM missing something simple...options include either replacing threadgroup matrices with our own custom type (like with safe arrays?), or adding a portability flag to indicate and test for workgroupMatrices support and direct CTS tests accordingly.

This hasn't come up for us with a real-world app AFAIK, so we might be able to get away with the latter.

[billhollings]

options include replacing threadgroup matrices with our own custom type (like with safe arrays?)

Referencing the thought here, another option might be to elevate local threadgroup structs to input/output variables passed into the Metal function as buffer arguments:

kernel void main1( threadgroup Foo& foo [[buffer(n)]] ) {