diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp new file mode 100644 index 00000000..5dd6113b --- /dev/null +++ b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp @@ -0,0 +1,52 @@ +#include +#include + +using namespace metal; + +struct A +{ + float v; + char _m0_final_padding[12]; +}; + +struct B +{ + float2 v; + char _m0_final_padding[8]; +}; + +struct C +{ + float3 v; +}; + +struct D +{ + float4 v; +}; + +struct E +{ + float4 a; + float2 b; + char _m0_final_padding[8]; +}; + +struct SSBO +{ + A a[2][4]; + B b[2][4]; + C c[2][4]; + D d[2][4]; + float2x4 e[2][4]; + E f[1]; +}; + +kernel void main0(device SSBO& _32 [[buffer(0)]]) +{ + _32.f[0].a = float4(2.0); + float2x2 tmp = float2x2(_32.e[0][1][0].xy, _32.e[0][1][1].xy); + _32.e[1][2][0].xy = tmp[0]; + _32.e[1][2][1].xy = tmp[1]; +} + diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp new file mode 100644 index 00000000..521ac018 --- /dev/null +++ b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp @@ -0,0 +1,52 @@ +#include +#include + +using namespace metal; + +struct A +{ + float v; + char _m0_final_padding[12]; +}; + +struct B +{ + float2 v; + char _m0_final_padding[8]; +}; + +struct C +{ + float3 v; +}; + +struct D +{ + float4 v; +}; + +struct E +{ + float4 a; + float2 b; + char _m0_final_padding[8]; +}; + +struct SSBO +{ + A a[4]; + B b[4]; + C c[4]; + D d[4]; + float2x4 e[4]; + E f[1]; +}; + +kernel void main0(device SSBO& _26 [[buffer(0)]]) +{ + _26.f[0].a = float4(2.0); + float2x2 tmp = float2x2(_26.e[1][0].xy, _26.e[1][1].xy); + _26.e[2][0].xy = tmp[0]; + _26.e[2][1].xy = tmp[1]; +} + diff --git a/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp b/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp new file mode 100644 index 00000000..ef1ba65c --- /dev/null +++ b/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp @@ -0,0 +1,45 @@ +#version 450 +layout(local_size_x = 1) in; + +struct A +{ + float v; +}; + +struct B +{ + vec2 v; +}; + +struct C +{ + vec3 v; +}; + +struct D +{ + vec4 v; +}; + +struct E +{ + vec4 a; + vec2 b; +}; + +layout(std140, set = 0, binding = 0) buffer SSBO +{ + A a[2][4]; + B b[2][4]; + C c[2][4]; + D d[2][4]; + mat2 e[2][4]; + E f[]; +}; + +void main() +{ + f[0].a = vec4(2.0); + mat2 tmp = e[0][1]; + e[1][2] = tmp; +} diff --git a/shaders-msl-no-opt/packing/struct-size-padding.comp b/shaders-msl-no-opt/packing/struct-size-padding.comp new file mode 100644 index 00000000..ad65415b --- /dev/null +++ b/shaders-msl-no-opt/packing/struct-size-padding.comp @@ -0,0 +1,45 @@ +#version 450 +layout(local_size_x = 1) in; + +struct A +{ + float v; +}; + +struct B +{ + vec2 v; +}; + +struct C +{ + vec3 v; +}; + +struct D +{ + vec4 v; +}; + +struct E +{ + vec4 a; + vec2 b; +}; + +layout(std140, set = 0, binding = 0) buffer SSBO +{ + A a[4]; + B b[4]; + C c[4]; + D d[4]; + mat2 e[4]; + E f[]; +}; + +void main() +{ + f[0].a = vec4(2.0); + mat2 tmp = e[1]; + e[2] = tmp; +} diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 41a67df3..6c67a7fc 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -6545,9 +6545,6 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice BuiltIn builtin; if (is_member_builtin(*type, index, &builtin)) { - // FIXME: We rely here on OpName on gl_in/gl_out to make this work properly. - // To make this properly work by omitting all OpName opcodes, - // we need to infer gl_in or gl_out based on the builtin, and stage. if (access_chain_is_arrayed) { expr += "."; @@ -6608,8 +6605,6 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice deferred_index = expr.substr(column_index); expr.resize(column_index); } - is_packed = false; - physical_type = 0; } if (index_is_literal && !is_packed && !row_major_matrix_needs_conversion) diff --git a/spirv_msl.cpp b/spirv_msl.cpp index c4dbca55..6af2db85 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -2427,10 +2427,10 @@ void CompilerMSL::mark_scalar_layout_structs(const SPIRType &type) if (!mbr_type.array.empty()) { array_stride = type_struct_member_array_stride(type, i); - uint32_t dimensions = type.array.size(); - for (uint32_t dim = 1; dim < dimensions; dim++) + uint32_t dimensions = mbr_type.array.size() - 1; + for (uint32_t dim = 0; dim < dimensions; dim++) { - uint32_t array_size = to_array_size_literal(type, dim); + uint32_t array_size = to_array_size_literal(mbr_type, dim); array_stride /= max(array_size, 1u); } @@ -2954,7 +2954,7 @@ string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type }; // std140 array cases for vectors. - if (physical_type && is_array(*physical_type) && physical_type->vecsize > type.vecsize) + if (physical_type && is_vector(*physical_type) && is_array(*physical_type) && physical_type->vecsize > type.vecsize) { assert(type.vecsize >= 1 && type.vecsize <= 3); return enclose_expression(expr_str) + swizzle_lut[type.vecsize - 1];