Skip to content

MSL: Rewrite std140 struct padding. #2598

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
HansKristian-Work merged 2 commits into from
Feb 12, 2026
+240 −155
Merged

MSL: Rewrite std140 struct padding. #2598

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
5ea5096
MSL: Rewrite std140 struct padding.
HansKristian-Work Feb 12, 2026
46c9d93
test: Fix shader_model_hlsl.
HansKristian-Work Feb 12, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 11 additions & 7 deletions reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,13 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct A
{
int a;
Expand All @@ -18,12 +23,11 @@ struct A_2
{
int a;
int b;
char _m0_final_padding[8];
};

struct A_3
{
A_2 Data[1024];
spvPaddedArrayElement<A_2, 16> Data[1024];
};

struct B
Expand All @@ -33,14 +37,14 @@ struct B

struct B_1
{
A_2 Data[1024];
spvPaddedArrayElement<A_2, 16> Data[1024];
};

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]])
{
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
}

36 changes: 18 additions & 18 deletions reference/opt/shaders-msl/comp/struct-packing.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,13 +1,17 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct S0
{
float2 a[1];
float b;
char _m0_final_padding[4];
};

struct S1
Expand All @@ -20,7 +24,6 @@ struct S2
{
float3 a[1];
float b;
char _m0_final_padding[12];
};

struct S3
Expand All @@ -45,7 +48,6 @@ struct Content
S3 m3;
float m4;
S4 m3s[8];
char _m0_final_padding[8];
};

struct SSBO1
Expand All @@ -69,7 +71,6 @@ struct S0_1
float2 a[1];
char _m1_pad[8];
float b;
char _m0_final_padding[12];
};

struct S1_1
Expand All @@ -82,7 +83,6 @@ struct S2_1
{
float3 a[1];
float b;
char _m0_final_padding[12];
};

struct S3_1
Expand All @@ -94,21 +94,21 @@ struct S3_1
struct S4_1
{
float2 c;
char _m0_final_padding[8];
};

struct Content_1
{
S0_1 m0s[1];
spvPaddedArrayElement<S0_1, 32> m0s[1];
S1_1 m1s[1];
S2_1 m2s[1];
S0_1 m0;
char _m4_pad[8];
S1_1 m1;
S2_1 m2;
S3_1 m3;
float m4;
char _m8_pad[8];
S4_1 m3s[8];
spvPaddedArrayElement<S4_1, 16> m3s[8];
};

struct SSBO0
Expand All @@ -124,8 +124,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
{
Content_1 _60 = ssbo_140.content;
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].data.a[0];
ssbo_430.content.m0s[0].b = _60.m0s[0].data.b;
ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
Expand All @@ -139,14 +139,14 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
ssbo_430.content.m3.a = _60.m3.a;
ssbo_430.content.m3.b = _60.m3.b;
ssbo_430.content.m4 = _60.m4;
ssbo_430.content.m3s[0].c = _60.m3s[0].c;
ssbo_430.content.m3s[1].c = _60.m3s[1].c;
ssbo_430.content.m3s[2].c = _60.m3s[2].c;
ssbo_430.content.m3s[3].c = _60.m3s[3].c;
ssbo_430.content.m3s[4].c = _60.m3s[4].c;
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
ssbo_430.content.m3s[0].c = _60.m3s[0].data.c;
ssbo_430.content.m3s[1].c = _60.m3s[1].data.c;
ssbo_430.content.m3s[2].c = _60.m3s[2].data.c;
ssbo_430.content.m3s[3].c = _60.m3s[3].data.c;
ssbo_430.content.m3s[4].c = _60.m3s[4].data.c;
ssbo_430.content.m3s[5].c = _60.m3s[5].data.c;
ssbo_430.content.m3s[6].c = _60.m3s[6].data.c;
ssbo_430.content.m3s[7].c = _60.m3s[7].data.c;
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
}

10 changes: 7 additions & 3 deletions reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,13 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct T
{
float a;
Expand All @@ -21,18 +26,17 @@ struct SSBO1
struct T_2
{
float c;
char _m0_final_padding[12];
};

struct SSBO2
{
T_2 bar[1];
spvPaddedArrayElement<T_2, 16> bar[1];
};

kernel void main0(device SSBO1& _9 [[buffer(0)]], device SSBO2& _13 [[buffer(1)]])
{
T v = T{ 40.0 };
_9.foo[10].b = v.a;
_13.bar[30].c = v.a;
_13.bar[30].data.c = v.a;
}

15 changes: 9 additions & 6 deletions reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,5 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

Expand All @@ -7,6 +9,9 @@ typedef packed_float3 packed_float2x3[2];
typedef packed_float3 packed_rm_float3x2[2];
typedef packed_float2 packed_float2x2[2];

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct S0
{
packed_float2 a[1];
Expand Down Expand Up @@ -64,7 +69,6 @@ struct S0_1
float2 a[1];
char _m1_pad[8];
float b;
char _m0_final_padding[12];
};

struct S1_1
Expand All @@ -77,7 +81,6 @@ struct S2_1
{
float3 a[1];
float b;
char _m0_final_padding[12];
};

struct S3_1
Expand All @@ -88,15 +91,15 @@ struct S3_1

struct Content_1
{
S0_1 m0s[1];
spvPaddedArrayElement<S0_1, 32> m0s[1];
S1_1 m1s[1];
S2_1 m2s[1];
S0_1 m0;
char _m4_pad[8];
S1_1 m1;
S2_1 m2;
S3_1 m3;
float m4;
char _m0_final_padding[12];
};

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

kernel void main0(device SSBO1& __restrict ssbo_scalar [[buffer(0)]], device SSBO0& __restrict ssbo_140 [[buffer(1)]], device SSBO2& __restrict ssbo_scalar2 [[buffer(2)]])
{
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].data.a[0];
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].data.b;
ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
Expand Down
12 changes: 7 additions & 5 deletions reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,18 +1,21 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct A
{
float v;
char _m0_final_padding[12];
};

struct B
{
float2 v;
char _m0_final_padding[8];
};

struct C
Expand All @@ -29,13 +32,12 @@ struct E
{
float4 a;
float2 b;
char _m0_final_padding[8];
};

struct SSBO
{
A a[2][4];
B b[2][4];
spvPaddedArrayElement<A, 16> a[2][4];
spvPaddedArrayElement<B, 16> b[2][4];
C c[2][4];
D d[2][4];
float2x4 e[2][4];
Expand Down
12 changes: 7 additions & 5 deletions reference/shaders-msl-no-opt/packing/struct-size-padding.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,18 +1,21 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct A
{
float v;
char _m0_final_padding[12];
};

struct B
{
float2 v;
char _m0_final_padding[8];
};

struct C
Expand All @@ -29,13 +32,12 @@ struct E
{
float4 a;
float2 b;
char _m0_final_padding[8];
};

struct SSBO
{
A a[4];
B b[4];
spvPaddedArrayElement<A, 16> a[4];
spvPaddedArrayElement<B, 16> b[4];
C c[4];
D d[4];
float2x4 e[4];
Expand Down
18 changes: 11 additions & 7 deletions reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,13 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template <typename T, int stride>
struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };

struct A
{
int a;
Expand All @@ -18,12 +23,11 @@ struct A_2
{
int a;
int b;
char _m0_final_padding[8];
};

struct A_3
{
A_2 Data[1024];
spvPaddedArrayElement<A_2, 16> Data[1024];
};

struct B
Expand All @@ -33,14 +37,14 @@ struct B

struct B_1
{
A_2 Data[1024];
spvPaddedArrayElement<A_2, 16> Data[1024];
};

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]])
{
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
}

Loading