Skip to content

Commit bf6bb5c

Browse files
HansKristian-WorkHansKristian-Work
authored
Merge pull request #2598 from KhronosGroup/fix-2595
MSL: Rewrite std140 struct padding.
2 parents f6e58e0 + 46c9d93 commit bf6bb5c

13 files changed

Lines changed: 240 additions & 155 deletions

‎reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp‎

Lines changed: 11 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,13 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct A
712
{
813
int a;
@@ -18,12 +23,11 @@ struct A_2
1823
{
1924
int a;
2025
int b;
21-
char _m0_final_padding[8];
2226
};
2327

2428
struct A_3
2529
{
26-
A_2 Data[1024];
30+
spvPaddedArrayElement<A_2, 16> Data[1024];
2731
};
2832

2933
struct B
@@ -33,14 +37,14 @@ struct B
3337

3438
struct B_1
3539
{
36-
A_2 Data[1024];
40+
spvPaddedArrayElement<A_2, 16> Data[1024];
3741
};
3842

3943
kernel void main0(device A_1& C1 [[buffer(0)]], constant A_3& C2 [[buffer(1)]], device B& C3 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
4044
{
41-
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
42-
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
43-
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
44-
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
45+
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
46+
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
47+
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
48+
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
4549
}
4650

‎reference/opt/shaders-msl/comp/struct-packing.comp‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,17 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct S0
712
{
813
float2 a[1];
914
float b;
10-
char _m0_final_padding[4];
1115
};
1216

1317
struct S1
@@ -20,7 +24,6 @@ struct S2
2024
{
2125
float3 a[1];
2226
float b;
23-
char _m0_final_padding[12];
2427
};
2528

2629
struct S3
@@ -45,7 +48,6 @@ struct Content
4548
S3 m3;
4649
float m4;
4750
S4 m3s[8];
48-
char _m0_final_padding[8];
4951
};
5052

5153
struct SSBO1
@@ -69,7 +71,6 @@ struct S0_1
6971
float2 a[1];
7072
char _m1_pad[8];
7173
float b;
72-
char _m0_final_padding[12];
7374
};
7475

7576
struct S1_1
@@ -82,7 +83,6 @@ struct S2_1
8283
{
8384
float3 a[1];
8485
float b;
85-
char _m0_final_padding[12];
8686
};
8787

8888
struct S3_1
@@ -94,21 +94,21 @@ struct S3_1
9494
struct S4_1
9595
{
9696
float2 c;
97-
char _m0_final_padding[8];
9897
};
9998

10099
struct Content_1
101100
{
102-
S0_1 m0s[1];
101+
spvPaddedArrayElement<S0_1, 32> m0s[1];
103102
S1_1 m1s[1];
104103
S2_1 m2s[1];
105104
S0_1 m0;
105+
char _m4_pad[8];
106106
S1_1 m1;
107107
S2_1 m2;
108108
S3_1 m3;
109109
float m4;
110110
char _m8_pad[8];
111-
S4_1 m3s[8];
111+
spvPaddedArrayElement<S4_1, 16> m3s[8];
112112
};
113113

114114
struct SSBO0
@@ -124,8 +124,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
124124
kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
125125
{
126126
Content_1 _60 = ssbo_140.content;
127-
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
128-
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
127+
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].data.a[0];
128+
ssbo_430.content.m0s[0].b = _60.m0s[0].data.b;
129129
ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
130130
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
131131
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
@@ -139,14 +139,14 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
139139
ssbo_430.content.m3.a = _60.m3.a;
140140
ssbo_430.content.m3.b = _60.m3.b;
141141
ssbo_430.content.m4 = _60.m4;
142-
ssbo_430.content.m3s[0].c = _60.m3s[0].c;
143-
ssbo_430.content.m3s[1].c = _60.m3s[1].c;
144-
ssbo_430.content.m3s[2].c = _60.m3s[2].c;
145-
ssbo_430.content.m3s[3].c = _60.m3s[3].c;
146-
ssbo_430.content.m3s[4].c = _60.m3s[4].c;
147-
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
148-
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
149-
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
142+
ssbo_430.content.m3s[0].c = _60.m3s[0].data.c;
143+
ssbo_430.content.m3s[1].c = _60.m3s[1].data.c;
144+
ssbo_430.content.m3s[2].c = _60.m3s[2].data.c;
145+
ssbo_430.content.m3s[3].c = _60.m3s[3].data.c;
146+
ssbo_430.content.m3s[4].c = _60.m3s[4].data.c;
147+
ssbo_430.content.m3s[5].c = _60.m3s[5].data.c;
148+
ssbo_430.content.m3s[6].c = _60.m3s[6].data.c;
149+
ssbo_430.content.m3s[7].c = _60.m3s[7].data.c;
150150
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
151151
}
152152

‎reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp‎

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,13 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct T
712
{
813
float a;
@@ -21,18 +26,17 @@ struct SSBO1
2126
struct T_2
2227
{
2328
float c;
24-
char _m0_final_padding[12];
2529
};
2630

2731
struct SSBO2
2832
{
29-
T_2 bar[1];
33+
spvPaddedArrayElement<T_2, 16> bar[1];
3034
};
3135

3236
kernel void main0(device SSBO1& _9 [[buffer(0)]], device SSBO2& _13 [[buffer(1)]])
3337
{
3438
T v = T{ 40.0 };
3539
_9.foo[10].b = v.a;
36-
_13.bar[30].c = v.a;
40+
_13.bar[30].data.c = v.a;
3741
}
3842

‎reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp‎

Lines changed: 9 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

@@ -7,6 +9,9 @@ typedef packed_float3 packed_float2x3[2];
79
typedef packed_float3 packed_rm_float3x2[2];
810
typedef packed_float2 packed_float2x2[2];
911

12+
template <typename T, int stride>
13+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
14+
1015
struct S0
1116
{
1217
packed_float2 a[1];
@@ -64,7 +69,6 @@ struct S0_1
6469
float2 a[1];
6570
char _m1_pad[8];
6671
float b;
67-
char _m0_final_padding[12];
6872
};
6973

7074
struct S1_1
@@ -77,7 +81,6 @@ struct S2_1
7781
{
7882
float3 a[1];
7983
float b;
80-
char _m0_final_padding[12];
8184
};
8285

8386
struct S3_1
@@ -88,15 +91,15 @@ struct S3_1
8891

8992
struct Content_1
9093
{
91-
S0_1 m0s[1];
94+
spvPaddedArrayElement<S0_1, 32> m0s[1];
9295
S1_1 m1s[1];
9396
S2_1 m2s[1];
9497
S0_1 m0;
98+
char _m4_pad[8];
9599
S1_1 m1;
96100
S2_1 m2;
97101
S3_1 m3;
98102
float m4;
99-
char _m0_final_padding[12];
100103
};
101104

102105
struct SSBO0
@@ -126,8 +129,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
126129

127130
kernel void main0(device SSBO1& __restrict ssbo_scalar [[buffer(0)]], device SSBO0& __restrict ssbo_140 [[buffer(1)]], device SSBO2& __restrict ssbo_scalar2 [[buffer(2)]])
128131
{
129-
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
130-
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
132+
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].data.a[0];
133+
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].data.b;
131134
ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
132135
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
133136
ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];

‎reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp‎

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,18 +1,21 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct A
712
{
813
float v;
9-
char _m0_final_padding[12];
1014
};
1115

1216
struct B
1317
{
1418
float2 v;
15-
char _m0_final_padding[8];
1619
};
1720

1821
struct C
@@ -29,13 +32,12 @@ struct E
2932
{
3033
float4 a;
3134
float2 b;
32-
char _m0_final_padding[8];
3335
};
3436

3537
struct SSBO
3638
{
37-
A a[2][4];
38-
B b[2][4];
39+
spvPaddedArrayElement<A, 16> a[2][4];
40+
spvPaddedArrayElement<B, 16> b[2][4];
3941
C c[2][4];
4042
D d[2][4];
4143
float2x4 e[2][4];

‎reference/shaders-msl-no-opt/packing/struct-size-padding.comp‎

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,18 +1,21 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct A
712
{
813
float v;
9-
char _m0_final_padding[12];
1014
};
1115

1216
struct B
1317
{
1418
float2 v;
15-
char _m0_final_padding[8];
1619
};
1720

1821
struct C
@@ -29,13 +32,12 @@ struct E
2932
{
3033
float4 a;
3134
float2 b;
32-
char _m0_final_padding[8];
3335
};
3436

3537
struct SSBO
3638
{
37-
A a[4];
38-
B b[4];
39+
spvPaddedArrayElement<A, 16> a[4];
40+
spvPaddedArrayElement<B, 16> b[4];
3941
C c[4];
4042
D d[4];
4143
float2x4 e[4];

‎reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp‎

Lines changed: 11 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,13 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
13
#include <metal_stdlib>
24
#include <simd/simd.h>
35

46
using namespace metal;
57

8+
template <typename T, int stride>
9+
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
10+
611
struct A
712
{
813
int a;
@@ -18,12 +23,11 @@ struct A_2
1823
{
1924
int a;
2025
int b;
21-
char _m0_final_padding[8];
2226
};
2327

2428
struct A_3
2529
{
26-
A_2 Data[1024];
30+
spvPaddedArrayElement<A_2, 16> Data[1024];
2731
};
2832

2933
struct B
@@ -33,14 +37,14 @@ struct B
3337

3438
struct B_1
3539
{
36-
A_2 Data[1024];
40+
spvPaddedArrayElement<A_2, 16> Data[1024];
3741
};
3842

3943
kernel void main0(device A_1& C1 [[buffer(0)]], constant A_3& C2 [[buffer(1)]], device B& C3 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
4044
{
41-
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
42-
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
43-
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
44-
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
45+
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
46+
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
47+
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
48+
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
4549
}
4650

0 commit comments

Comments
 (0)