AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #1331 from KhronosGroup/fix-1321

Browse Source 
MSL: Support edge case with DX layout in scalar block layout.
This commit is contained in:
Hans-Kristian Arntzen 2020-04-21 10:50:47 +02:00 committed by GitHub
commit 3fb86e4385
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
26 changed files with 970 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
reference/opt/shaders-msl/comp/struct-packing.comp
View File

@ -66,7 +66,8 @@ struct SSBO1
struct S0_1
{
float4 a[1];
packed_float2 a[1];
char _m1_pad[8];
float b;
char _m0_final_padding[12];
};
@ -123,13 +124,13 @@ 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].xy;
ssbo_430.content.m0s[0].a[0] = float2(_60.m0s[0].a[0]);
ssbo_430.content.m0s[0].b = _60.m0s[0].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];
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
ssbo_430.content.m0.a[0] = _60.m0.a[0].xy;
ssbo_430.content.m0.a[0] = float2(_60.m0.a[0]);
ssbo_430.content.m0.b = _60.m0.b;
ssbo_430.content.m1.a = float3(_60.m1.a);
ssbo_430.content.m1.b = _60.m1.b;

24
reference/shaders-msl-no-opt/asm/packing/scalar-array-float2.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[8];
float b;
};
struct main0_out
{
float2 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0].xy + Foo.a[1].xy) + float2(Foo.b);
return out;
}

23
reference/shaders-msl-no-opt/asm/packing/scalar-array-float3-one-element.asm.frag Normal file
View File

@ -0,0 +1,23 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
packed_float3 a[1];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = float3(Foo.a[0]) + float3(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-array-float3.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[12];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0].xyz + Foo.a[1].xyz) + float3(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float2x2-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[8];
float b;
};
struct main0_out
{
float2 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0u].xy + Foo.a[1u].xy) + float2(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float2x2-row-major.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[8];
float b;
};
struct main0_out
{
float2 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (float2(Foo.a[0][0u], Foo.a[1][0u]) + float2(Foo.a[0][1u], Foo.a[1][1u])) + float2(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float2x3-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[12];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0u].xyz + Foo.a[1u].xyz) + float3(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float2x3-row-major.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float2x4 a;
char _m1_pad[8];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (float3(Foo.a[0][0u], Foo.a[1][0u], Foo.a[2][0u]) + float3(Foo.a[0][1u], Foo.a[1][1u], Foo.a[2][1u])) + float3(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float3x2-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float2x4 a;
char _m1_pad[8];
float b;
};
struct main0_out
{
float2 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0u].xy + Foo.a[1u].xy) + float2(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float3x2-row-major.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float4 a[1];
char _m1_pad[12];
float b;
};
struct main0_out
{
float2 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (float2(Foo.a[0][0u], Foo.a[1][0u]) + float2(Foo.a[0][1u], Foo.a[1][1u])) + float2(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float3x3-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float2x4 a;
char _m1_pad[12];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (Foo.a[0u].xyz + Foo.a[1u].xyz) + float3(Foo.b);
return out;
}

24
reference/shaders-msl-no-opt/asm/packing/scalar-float3x3-row-major.asm.frag Normal file
View File

@ -0,0 +1,24 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct type_Foo
{
float2x4 a;
char _m1_pad[12];
float b;
};
struct main0_out
{
float3 out_var_SV_Target [[color(0)]];
};
fragment main0_out main0(constant type_Foo& Foo [[buffer(0)]])
{
main0_out out = {};
out.out_var_SV_Target = (float3(Foo.a[0][0u], Foo.a[1][0u], Foo.a[2][0u]) + float3(Foo.a[0][1u], Foo.a[1][1u], Foo.a[2][1u])) + float3(Foo.b);
return out;
}

7
reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
View File

@ -61,7 +61,8 @@ struct SSBO1
struct S0_1
{
float4 a[1];
packed_float2 a[1];
char _m1_pad[8];
float b;
char _m0_final_padding[12];
};
@ -125,13 +126,13 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
{
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy;
ssbo_scalar.content.m0s[0].a[0] = float2(ssbo_140.content.m0s[0].a[0]);
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].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];
ssbo_scalar.content.m2s[0].b = ssbo_140.content.m2s[0].b;
ssbo_scalar.content.m0.a[0] = ssbo_140.content.m0.a[0].xy;
ssbo_scalar.content.m0.a[0] = float2(ssbo_140.content.m0.a[0]);
ssbo_scalar.content.m0.b = ssbo_140.content.m0.b;
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;

7
reference/shaders-msl/comp/struct-packing.comp
View File

@ -66,7 +66,8 @@ struct SSBO1
struct S0_1
{
float4 a[1];
packed_float2 a[1];
char _m1_pad[8];
float b;
char _m0_final_padding[12];
};
@ -123,13 +124,13 @@ 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].xy;
ssbo_430.content.m0s[0].a[0] = float2(_60.m0s[0].a[0]);
ssbo_430.content.m0s[0].b = _60.m0s[0].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];
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
ssbo_430.content.m0.a[0] = _60.m0.a[0].xy;
ssbo_430.content.m0.a[0] = float2(_60.m0.a[0]);
ssbo_430.content.m0.b = _60.m0.b;
ssbo_430.content.m1.a = float3(_60.m1.a);
ssbo_430.content.m1.b = _60.m1.b;

54
shaders-msl-no-opt/asm/packing/scalar-array-float2.asm.frag Normal file
View File

@ -0,0 +1,54 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 29
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpDecorate %_arr_v2float_uint_2 ArrayStride 16
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 1 Offset 24
OpDecorate %type_Foo Block
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%int_1 = OpConstant %int 1
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%_arr_v2float_uint_2 = OpTypeArray %v2float %uint_2
%type_Foo = OpTypeStruct %_arr_v2float_uint_2 %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v2float = OpTypePointer Output %v2float
%void = OpTypeVoid
%16 = OpTypeFunction %void
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v2float Output
%main = OpFunction %void None %16
%19 = OpLabel
%20 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %int_0
%21 = OpLoad %v2float %20
%22 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %int_1
%23 = OpLoad %v2float %22
%24 = OpFAdd %v2float %21 %23
%25 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%26 = OpLoad %float %25
%27 = OpCompositeConstruct %v2float %26 %26
%28 = OpFAdd %v2float %24 %27
OpStore %out_var_SV_Target %28
OpReturn
OpFunctionEnd

51
shaders-msl-no-opt/asm/packing/scalar-array-float3-one-element.asm.frag Normal file
View File

@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 26
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpDecorate %_arr_v3float_uint_1 ArrayStride 16
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 1 Offset 12
OpDecorate %type_Foo Block
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%int_1 = OpConstant %int 1
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%_arr_v3float_uint_1 = OpTypeArray %v3float %uint_1
%type_Foo = OpTypeStruct %_arr_v3float_uint_1 %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%16 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %16
%19 = OpLabel
%20 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %int_0
%21 = OpLoad %v3float %20
%22 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%23 = OpLoad %float %22
%24 = OpCompositeConstruct %v3float %23 %23 %23
%25 = OpFAdd %v3float %21 %24
OpStore %out_var_SV_Target %25
OpReturn
OpFunctionEnd

54
shaders-msl-no-opt/asm/packing/scalar-array-float3.asm.frag Normal file
View File

@ -0,0 +1,54 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 29
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpDecorate %_arr_v3float_uint_2 ArrayStride 16
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 1 Offset 28
OpDecorate %type_Foo Block
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%int_1 = OpConstant %int 1
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%_arr_v3float_uint_2 = OpTypeArray %v3float %uint_2
%type_Foo = OpTypeStruct %_arr_v3float_uint_2 %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%16 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %16
%19 = OpLabel
%20 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %int_0
%21 = OpLoad %v3float %20
%22 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %int_1
%23 = OpLoad %v3float %22
%24 = OpFAdd %v3float %21 %23
%25 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%26 = OpLoad %float %25
%27 = OpCompositeConstruct %v3float %26 %26 %26
%28 = OpFAdd %v3float %24 %27
OpStore %out_var_SV_Target %28
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float2x2-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 ColMajor
OpMemberDecorate %type_Foo 1 Offset 24
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%type_Foo = OpTypeStruct %mat2v2float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v2float = OpTypePointer Output %v2float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v2float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_0
%22 = OpLoad %v2float %21
%23 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_1
%24 = OpLoad %v2float %23
%25 = OpFAdd %v2float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v2float %27 %27
%29 = OpFAdd %v2float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float2x2-row-major.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 RowMajor
OpMemberDecorate %type_Foo 1 Offset 24
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%type_Foo = OpTypeStruct %mat2v2float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v2float = OpTypePointer Output %v2float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v2float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_0
%22 = OpLoad %v2float %21
%23 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_1
%24 = OpLoad %v2float %23
%25 = OpFAdd %v2float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v2float %27 %27
%29 = OpFAdd %v2float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float2x3-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 ColMajor
OpMemberDecorate %type_Foo 1 Offset 28
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%mat2v3float = OpTypeMatrix %v3float 2
%type_Foo = OpTypeStruct %mat2v3float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_0
%22 = OpLoad %v3float %21
%23 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_1
%24 = OpLoad %v3float %23
%25 = OpFAdd %v3float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v3float %27 %27 %27
%29 = OpFAdd %v3float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float2x3-row-major.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 RowMajor
OpMemberDecorate %type_Foo 1 Offset 40
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%mat2v3float = OpTypeMatrix %v3float 2
%type_Foo = OpTypeStruct %mat2v3float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_0
%22 = OpLoad %v3float %21
%23 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_1
%24 = OpLoad %v3float %23
%25 = OpFAdd %v3float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v3float %27 %27 %27
%29 = OpFAdd %v3float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float3x2-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 ColMajor
OpMemberDecorate %type_Foo 1 Offset 40
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat3v2float = OpTypeMatrix %v2float 3
%type_Foo = OpTypeStruct %mat3v2float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v2float = OpTypePointer Output %v2float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v2float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_0
%22 = OpLoad %v2float %21
%23 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_1
%24 = OpLoad %v2float %23
%25 = OpFAdd %v2float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v2float %27 %27
%29 = OpFAdd %v2float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float3x2-row-major.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 RowMajor
OpMemberDecorate %type_Foo 1 Offset 28
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat3v2float = OpTypeMatrix %v2float 3
%type_Foo = OpTypeStruct %mat3v2float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v2float = OpTypePointer Output %v2float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v2float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_0
%22 = OpLoad %v2float %21
%23 = OpAccessChain %_ptr_Uniform_v2float %Foo %int_0 %uint_1
%24 = OpLoad %v2float %23
%25 = OpFAdd %v2float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v2float %27 %27
%29 = OpFAdd %v2float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float3x3-col-major.invalid.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 ColMajor
OpMemberDecorate %type_Foo 1 Offset 44
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%mat3v3float = OpTypeMatrix %v3float 3
%type_Foo = OpTypeStruct %mat3v3float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_0
%22 = OpLoad %v3float %21
%23 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_1
%24 = OpLoad %v3float %23
%25 = OpFAdd %v3float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v3float %27 %27 %27
%29 = OpFAdd %v3float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

56
shaders-msl-no-opt/asm/packing/scalar-float3x3-row-major.asm.frag Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.0
; Generator: Google spiregg; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %out_var_SV_Target
OpExecutionMode %main OriginUpperLeft
OpSource HLSL 600
OpName %type_Foo "type.Foo"
OpMemberName %type_Foo 0 "a"
OpMemberName %type_Foo 1 "b"
OpName %Foo "Foo"
OpName %out_var_SV_Target "out.var.SV_Target"
OpName %main "main"
OpDecorate %out_var_SV_Target Location 0
OpDecorate %Foo DescriptorSet 0
OpDecorate %Foo Binding 0
OpMemberDecorate %type_Foo 0 Offset 0
OpMemberDecorate %type_Foo 0 MatrixStride 16
OpMemberDecorate %type_Foo 0 RowMajor
OpMemberDecorate %type_Foo 1 Offset 44
OpDecorate %type_Foo Block
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%int_1 = OpConstant %int 1
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%mat3v3float = OpTypeMatrix %v3float 3
%type_Foo = OpTypeStruct %mat3v3float %float
%_ptr_Uniform_type_Foo = OpTypePointer Uniform %type_Foo
%_ptr_Output_v3float = OpTypePointer Output %v3float
%void = OpTypeVoid
%17 = OpTypeFunction %void
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%Foo = OpVariable %_ptr_Uniform_type_Foo Uniform
%out_var_SV_Target = OpVariable %_ptr_Output_v3float Output
%main = OpFunction %void None %17
%20 = OpLabel
%21 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_0
%22 = OpLoad %v3float %21
%23 = OpAccessChain %_ptr_Uniform_v3float %Foo %int_0 %uint_1
%24 = OpLoad %v3float %23
%25 = OpFAdd %v3float %22 %24
%26 = OpAccessChain %_ptr_Uniform_float %Foo %int_1
%27 = OpLoad %float %26
%28 = OpCompositeConstruct %v3float %27 %27 %27
%29 = OpFAdd %v3float %25 %28
OpStore %out_var_SV_Target %29
OpReturn
OpFunctionEnd

95
spirv_msl.cpp
View File

@ -3086,10 +3086,20 @@ bool CompilerMSL::validate_member_packing_rules_msl(const SPIRType &type, uint32
if (!mbr_type.array.empty())
{
// If we have an array type, array stride must match exactly with SPIR-V.
uint32_t spirv_array_stride = type_struct_member_array_stride(type, index);
uint32_t msl_array_stride = get_declared_struct_member_array_stride_msl(type, index);
if (spirv_array_stride != msl_array_stride)
return false;
// An exception to this requirement is if we have one array element and a packed decoration.
// This comes from DX scalar layout workaround.
// If app tries to be cheeky and access the member out of bounds, this will not work, but this is the best we can do.
bool relax_array_stride = has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPhysicalTypePacked) &&
mbr_type.array.back() == 1 && mbr_type.array_size_literal.back();
if (!relax_array_stride)
{
uint32_t spirv_array_stride = type_struct_member_array_stride(type, index);
uint32_t msl_array_stride = get_declared_struct_member_array_stride_msl(type, index);
if (spirv_array_stride != msl_array_stride)
return false;
}
}
if (is_matrix(mbr_type))
@ -3201,6 +3211,77 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
else
unset_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypePacked);
}
else
SPIRV_CROSS_THROW("Found a buffer packing case which we cannot represent in MSL.");
// Try validating again, now with physical type remapping.
if (validate_member_packing_rules_msl(ib_type, index))
return;
// We might have a particular odd scalar layout case where the last element of an array
// does not take up as much space as the ArrayStride or MatrixStride. This can happen with DX cbuffers.
// The "proper" workaround for this is extremely painful and essentially impossible in the edge case of float3[],
// so we hack around it by declaring the offending array or matrix with one less array size/col/row,
// and rely on padding to get the correct value. We will technically access arrays out of bounds into the padding region,
// but it should spill over gracefully without too much trouble. We rely on behavior like this for unsized arrays anyways.
// E.g. we might observe a physical layout of:
// { float2 a[2]; float b; } in cbuffer layout where ArrayStride of a is 16, but offset of b is 24, packed right after a[1] ...
uint32_t type_id = get_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypeID);
auto &type = get<SPIRType>(type_id);
// Modify the physical type in-place. This is safe since each physical type workaround is a copy.
if (is_array(type))
{
if (type.array.back() > 1)
{
if (!type.array_size_literal.back())
SPIRV_CROSS_THROW("Cannot apply scalar layout workaround with spec constant array size.");
type.array.back() -= 1;
}
else
{
// We have an array of size 1, so we cannot decrement that. Our only option now is to
// force a packed layout instead, and drop the physical type remap since ArrayStride is meaningless now.
unset_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypeID);
set_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPhysicalTypePacked);
}
}
else if (is_matrix(type))
{
bool row_major = has_member_decoration(ib_type.self, index, DecorationRowMajor);
if (!row_major)
{
// Slice off one column. If we only have 2 columns, this might turn the matrix into a vector with one array element instead.
if (type.columns > 2)
{
type.columns--;
}
else if (type.columns == 2)
{
type.columns = 1;
assert(type.array.empty());
type.array.push_back(1);
type.array_size_literal.push_back(true);
}
}
else
{
// Slice off one row. If we only have 2 rows, this might turn the matrix into a vector with one array element instead.
if (type.vecsize > 2)
{
type.vecsize--;
}
else if (type.vecsize == 2)
{
type.vecsize = type.columns;
type.columns = 1;
assert(type.array.empty());
type.array.push_back(1);
type.array_size_literal.push_back(true);
}
}
}
// This better validate now, or we must fail gracefully.
if (!validate_member_packing_rules_msl(ib_type, index))
@ -11824,7 +11905,7 @@ uint32_t CompilerMSL::get_declared_type_matrix_stride_msl(const SPIRType &type,
// For packed matrices, we just use the size of the vector type.
// Otherwise, MatrixStride == alignment, which is the size of the underlying vector type.
if (packed)
return (type.width / 8) * (row_major ? type.columns : type.vecsize);
return (type.width / 8) * ((row_major && type.columns > 1) ? type.columns : type.vecsize);
else
return get_declared_type_alignment_msl(type, false, row_major);
}
@ -11902,7 +11983,7 @@ uint32_t CompilerMSL::get_declared_type_size_msl(const SPIRType &type, bool is_p
uint32_t vecsize = type.vecsize;
uint32_t columns = type.columns;
if (row_major)
if (row_major && columns > 1)
swap(vecsize, columns);
if (vecsize == 3)
@ -11963,7 +12044,7 @@ uint32_t CompilerMSL::get_declared_type_alignment_msl(const SPIRType &type, bool
else
{
// This is the general rule for MSL. Size == alignment.
uint32_t vecsize = row_major ? type.columns : type.vecsize;
uint32_t vecsize = (row_major && type.columns > 1) ? type.columns : type.vecsize;
return (type.width / 8) * (vecsize == 3 ? 4 : vecsize);
}
}