diff --git a/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp index 95f2717b..e983acd9 100644 --- a/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp +++ b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp @@ -22,7 +22,7 @@ struct A_2 struct A_3 { - A_2 Data[1024]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024]; }; struct B @@ -32,7 +32,7 @@ struct B struct B_1 { - A_2 Data[1024]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024]; }; kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) diff --git a/reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag b/reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag index d7d6ef63..989d8268 100644 --- a/reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag +++ b/reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag @@ -73,7 +73,7 @@ struct _18 float2 _m24; float2 _m25; float2 _m26; - char pad27[8]; + char _m27_pad[8]; packed_float3 _m27; float _m28; float _m29; diff --git a/reference/opt/shaders-msl/comp/struct-packing.comp b/reference/opt/shaders-msl/comp/struct-packing.comp index 468eb7e6..bfaf4dd5 100644 --- a/reference/opt/shaders-msl/comp/struct-packing.comp +++ b/reference/opt/shaders-msl/comp/struct-packing.comp @@ -43,7 +43,6 @@ struct Content S1 m1; S2 m2; S3 m3; - char pad7[4]; float m4; S4 m3s[8]; }; @@ -53,7 +52,6 @@ struct SSBO1 Content content; Content content1[2]; Content content2; - char pad3[8]; float2x2 m0; float2x2 m1; float2x3 m2[4]; @@ -61,9 +59,9 @@ struct SSBO1 float2x2 m4; float2x2 m5[9]; packed_float2x3 m6[4][2]; - char pad10[8]; + char _m10_pad[8]; float3x2 m7; - char pad11[8]; + char _m11_pad[8]; float array[1]; }; @@ -105,9 +103,9 @@ struct Content_1 S1_1 m1; S2_1 m2; S3_1 m3; - char pad7[4]; float m4; - S4_1 m3s[8]; + char _m8_pad[12]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8]; }; struct SSBO0 diff --git a/reference/opt/shaders-msl/vert/packed_matrix.vert b/reference/opt/shaders-msl/vert/packed_matrix.vert index 7f058061..d319839c 100644 --- a/reference/opt/shaders-msl/vert/packed_matrix.vert +++ b/reference/opt/shaders-msl/vert/packed_matrix.vert @@ -16,7 +16,7 @@ struct _42 float4x4 _m0; float4x4 _m1; float _m2; - char pad3[12]; + char _m3_pad[12]; packed_float3 _m3; float _m4; packed_float3 _m5; diff --git a/reference/opt/shaders-msl/vert/ubo.alignment.vert b/reference/opt/shaders-msl/vert/ubo.alignment.vert index 9a7ea56c..c48111ed 100644 --- a/reference/opt/shaders-msl/vert/ubo.alignment.vert +++ b/reference/opt/shaders-msl/vert/ubo.alignment.vert @@ -7,7 +7,7 @@ struct UBO { float4x4 mvp; float2 targSize; - char pad2[8]; + char _m2_pad[8]; packed_float3 color; float opacity; }; diff --git a/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp index 95f2717b..e983acd9 100644 --- a/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp +++ b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp @@ -22,7 +22,7 @@ struct A_2 struct A_3 { - A_2 Data[1024]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024]; }; struct B @@ -32,7 +32,7 @@ struct B struct B_1 { - A_2 Data[1024]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ A_2 Data[1024]; }; kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) diff --git a/reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag b/reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag index 74bd618c..1738478d 100644 --- a/reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag +++ b/reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag @@ -73,7 +73,7 @@ struct _18 float2 _m24; float2 _m25; float2 _m26; - char pad27[8]; + char _m27_pad[8]; packed_float3 _m27; float _m28; float _m29; diff --git a/reference/shaders-msl/comp/struct-packing.comp b/reference/shaders-msl/comp/struct-packing.comp index 468eb7e6..bfaf4dd5 100644 --- a/reference/shaders-msl/comp/struct-packing.comp +++ b/reference/shaders-msl/comp/struct-packing.comp @@ -43,7 +43,6 @@ struct Content S1 m1; S2 m2; S3 m3; - char pad7[4]; float m4; S4 m3s[8]; }; @@ -53,7 +52,6 @@ struct SSBO1 Content content; Content content1[2]; Content content2; - char pad3[8]; float2x2 m0; float2x2 m1; float2x3 m2[4]; @@ -61,9 +59,9 @@ struct SSBO1 float2x2 m4; float2x2 m5[9]; packed_float2x3 m6[4][2]; - char pad10[8]; + char _m10_pad[8]; float3x2 m7; - char pad11[8]; + char _m11_pad[8]; float array[1]; }; @@ -105,9 +103,9 @@ struct Content_1 S1_1 m1; S2_1 m2; S3_1 m3; - char pad7[4]; float m4; - S4_1 m3s[8]; + char _m8_pad[12]; + /* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8]; }; struct SSBO0 diff --git a/reference/shaders-msl/vert/packed_matrix.vert b/reference/shaders-msl/vert/packed_matrix.vert index 89638511..98d4de53 100644 --- a/reference/shaders-msl/vert/packed_matrix.vert +++ b/reference/shaders-msl/vert/packed_matrix.vert @@ -16,7 +16,7 @@ struct _42 float4x4 _m0; float4x4 _m1; float _m2; - char pad3[12]; + char _m3_pad[12]; packed_float3 _m3; float _m4; packed_float3 _m5; diff --git a/reference/shaders-msl/vert/ubo.alignment.vert b/reference/shaders-msl/vert/ubo.alignment.vert index 9a7ea56c..c48111ed 100644 --- a/reference/shaders-msl/vert/ubo.alignment.vert +++ b/reference/shaders-msl/vert/ubo.alignment.vert @@ -7,7 +7,7 @@ struct UBO { float4x4 mvp; float2 targSize; - char pad2[8]; + char _m2_pad[8]; packed_float3 color; float opacity; }; diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 131bc803..e4d54e0d 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -960,7 +960,7 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits if (type.basetype == SPIRType::Struct) { // Rule 9. Structs alignments are maximum alignment of its members. - uint32_t alignment = 0; + uint32_t alignment = 1; for (uint32_t i = 0; i < type.member_types.size(); i++) { auto member_flags = ir.meta[type.self].members[i].decoration_flags; diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 923db5b1..7f2d3c6f 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -1565,11 +1565,11 @@ void CompilerMSL::align_struct(SPIRType &ib_type) // Align current offset to the current member's default alignment. size_t align_mask = get_declared_struct_member_alignment(ib_type, mbr_idx) - 1; - curr_offset = uint32_t((curr_offset + align_mask) & ~align_mask); + uint32_t aligned_curr_offset = uint32_t((curr_offset + align_mask) & ~align_mask); // Fetch the member offset as declared in the SPIRV. uint32_t mbr_offset = get_member_decoration(ib_type_id, mbr_idx, DecorationOffset); - if (mbr_offset > curr_offset) + if (mbr_offset > aligned_curr_offset) { // Since MSL and SPIR-V have slightly different struct member alignment and // size rules, we'll pad to standard C-packing rules. If the member is farther @@ -1610,6 +1610,18 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index) return true; } + // Check for array of struct, where the SPIR-V declares an array stride which is larger than the struct itself. + // This can happen for struct A { float a }; A a[]; in std140 layout. + // TODO: Emit a padded struct which can be used for this purpose. + if (is_array(mbr_type) && mbr_type.basetype == SPIRType::Struct) + { + size_t declared_struct_size = get_declared_struct_size(mbr_type); + size_t alignment = get_declared_struct_member_alignment(ib_type, index); + declared_struct_size = (declared_struct_size + alignment - 1) & ~(alignment - 1); + if (type_struct_member_array_stride(ib_type, index) > declared_struct_size) + return true; + } + // TODO: Another sanity check for matrices. We currently do not support std140 matrices which need to be padded out per column. //if (is_matrix(mbr_type) && mbr_type.vecsize <= 2 && type_struct_member_matrix_stride(ib_type, index) == 16) // SPIRV_CROSS_THROW("Currently cannot support matrices with small vector size in std140 layout."); @@ -4030,7 +4042,7 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_ MSLStructMemberKey key = get_struct_member_key(type.self, index); uint32_t pad_len = struct_member_padding[key]; if (pad_len > 0) - statement("char pad", to_string(index), "[", to_string(pad_len), "];"); + statement("char _m", index, "_pad", "[", to_string(pad_len), "];"); // If this member is packed, mark it as so. string pack_pfx = ""; @@ -4041,7 +4053,11 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_ if (member_is_packed_type(type, index)) { // If we're packing a matrix, output an appropriate typedef - if (membertype.vecsize > 1 && membertype.columns > 1) + if (membertype.basetype == SPIRType::Struct) + { + pack_pfx = "/* FIXME: A padded struct is needed here. If you see this message, file a bug! */ "; + } + else if (membertype.vecsize > 1 && membertype.columns > 1) { pack_pfx = "packed_"; string base_type = membertype.width == 16 ? "half" : "float"; @@ -5681,7 +5697,12 @@ size_t CompilerMSL::get_declared_struct_member_size(const SPIRType &struct_type, } if (type.basetype == SPIRType::Struct) - return get_declared_struct_size(type); + { + // The size of a struct in Metal is aligned up to its natural alignment. + auto size = get_declared_struct_size(type); + auto alignment = get_declared_struct_member_alignment(struct_type, index); + return (size + alignment - 1) & ~(alignment - 1); + } uint32_t component_size = type.width / 8; uint32_t vecsize = type.vecsize; @@ -5712,7 +5733,13 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_ SPIRV_CROSS_THROW("Querying alignment of opaque object."); case SPIRType::Struct: - return 16; // Per Vulkan spec section 14.5.4 + { + // In MSL, a struct's alignment is equal to the maximum alignment of any of its members. + uint32_t alignment = 1; + for (uint32_t i = 0; i < type.member_types.size(); i++) + alignment = max(alignment, uint32_t(get_declared_struct_member_alignment(type, i))); + return alignment; + } default: { @@ -5720,7 +5747,18 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_ // Alignment of unpacked type is the same as the vector size. // Alignment of 3-elements vector is the same as 4-elements (including packed using column). if (member_is_packed_type(struct_type, index)) - return (type.width / 8) * (type.columns == 3 ? 4 : type.columns); + { + // This is getting pretty complicated. + // The special case of array of float/float2 needs to be handled here. + uint32_t packed_type_id = + get_extended_member_decoration(struct_type.self, index, SPIRVCrossDecorationPackedType); + const SPIRType *packed_type = packed_type_id != 0 ? &get(packed_type_id) : nullptr; + if (packed_type && is_array(*packed_type) && !is_matrix(*packed_type) && + packed_type->basetype != SPIRType::Struct) + return (packed_type->width / 8) * 4; + else + return (type.width / 8) * (type.columns == 3 ? 4 : type.columns); + } else return (type.width / 8) * (type.vecsize == 3 ? 4 : type.vecsize); }