MSL: Support edge case with DX layout in scalar block layout.
DX may emit ArrayStride and MatrixStride of 16, but the size of the object does not align with that and expect to pack other members inside its last member. The workaround is to emit array size/col/row one less than we expect and rely on padding to carve out a "dead zone" for the last member.
This commit is contained in:
Hans-Kristian Arntzen
parent
f38cbeb814
commit
17ad62eea4
@ -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;
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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;
|
||||
|
||||
@ -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
54
shaders-msl-no-opt/asm/packing/scalar-array-float2.asm.frag
Normal 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
|
||||
@ -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
54
shaders-msl-no-opt/asm/packing/scalar-array-float3.asm.frag
Normal 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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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
|
||||
@ -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);
|
||||
}
|
||||
}
|
||||
|
||||
Loading…
Reference in New Issue
Block a user