--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float4 m_location_0 [[color(0)]];
+ float4 m_location_1 [[color(1)]];
+ float4 m_location_2 [[color(2)]];
+};
+
+struct main0_in
+{
+ float InC_0 [[user(locn0_1), flat]];
+ float InA_0 [[user(locn1), flat]];
+ float InC_1 [[user(locn1_1), flat]];
+ float2 InB_0 [[user(locn1_2), flat]];
+ float InA_1 [[user(locn2), flat]];
+ float InC_2 [[user(locn2_1), flat]];
+ float2 InB_1 [[user(locn2_2), flat]];
+ float InD [[user(locn3_1), sample_perspective]];
+ float InE [[user(locn4_2), center_no_perspective]];
+ float InF [[user(locn5_3), centroid_perspective]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+ main0_out out = {};
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ float D = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ InA[0] = in.InA_0;
+ InA[1] = in.InA_1;
+ InB[0] = in.InB_0;
+ InB[1] = in.InB_1;
+ InC[0] = in.InC_0;
+ InC[1] = in.InC_1;
+ InC[2] = in.InC_2;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = (in.InD + in.InE) + in.InF;
+ out.m_location_1.x = A[0];
+ out.m_location_2.x = A[1];
+ out.m_location_1.zw = B[0];
+ out.m_location_2.zw = B[1];
+ out.m_location_0.y = C[0];
+ out.m_location_1.y = C[1];
+ out.m_location_2.y = C[2];
+ out.m_location_0.w = D;
+ return out;
+}
+
--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float4 m_location_0;
+ float4 m_location_1;
+ float4 m_location_2;
+ float4 gl_Position;
+};
+
+struct main0_in
+{
+ float4 m_location_0 [[attribute(0)]];
+ float4 m_location_1 [[attribute(1)]];
+ float4 m_location_2 [[attribute(2)]];
+ float4 Pos [[attribute(4)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
+{
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ float D = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ float InD = {};
+ device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
+ InA[0] = in.m_location_1.x;
+ InA[1] = in.m_location_2.x;
+ InB[0] = in.m_location_1.zw;
+ InB[1] = in.m_location_2.zw;
+ InC[0] = in.m_location_0.y;
+ InC[1] = in.m_location_1.y;
+ InC[2] = in.m_location_2.y;
+ InD = in.m_location_0.w;
+ if (any(gl_GlobalInvocationID >= spvStageInputSize))
+ return;
+ out.gl_Position = in.Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = InD;
+ out.m_location_1.x = A[0];
+ out.m_location_2.x = A[1];
+ out.m_location_1.zw = B[0];
+ out.m_location_2.zw = B[1];
+ out.m_location_0.y = C[0];
+ out.m_location_1.y = C[1];
+ out.m_location_2.y = C[2];
+ out.m_location_0.w = D;
+}
+
--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float C_0 [[user(locn0_1)]];
+ float D [[user(locn0_3)]];
+ float A_0 [[user(locn1)]];
+ float C_1 [[user(locn1_1)]];
+ float2 B_0 [[user(locn1_2)]];
+ float A_1 [[user(locn2)]];
+ float C_2 [[user(locn2_1)]];
+ float2 B_1 [[user(locn2_2)]];
+ float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+ float4 m_location_0 [[attribute(0)]];
+ float4 m_location_1 [[attribute(1)]];
+ float4 m_location_2 [[attribute(2)]];
+ float4 Pos [[attribute(4)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+ main0_out out = {};
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ float InD = {};
+ InA[0] = in.m_location_1.x;
+ InA[1] = in.m_location_2.x;
+ InB[0] = in.m_location_1.zw;
+ InB[1] = in.m_location_2.zw;
+ InC[0] = in.m_location_0.y;
+ InC[1] = in.m_location_1.y;
+ InC[2] = in.m_location_2.y;
+ InD = in.m_location_0.w;
+ out.gl_Position = in.Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ out.D = InD;
+ out.A_0 = A[0];
+ out.A_1 = A[1];
+ out.B_0 = B[0];
+ out.B_1 = B[1];
+ out.C_0 = C[0];
+ out.C_1 = C[1];
+ out.C_2 = C[2];
+ return out;
+}
+
struct main0_out
{
- float4 FragColor0 [[color(0)]];
+ float4 m_location_0 [[color(0)]];
};
fragment main0_out main0()
FragColor0 = 1.0;
FragColor1 = float2(2.0, 3.0);
FragColor3 = 4.0;
- out.FragColor0.x = FragColor0;
- out.FragColor0.yz = FragColor1;
- out.FragColor0.w = FragColor3;
+ out.m_location_0.x = FragColor0;
+ out.m_location_0.yz = FragColor1;
+ out.m_location_0.w = FragColor3;
return out;
}
struct main0_out
{
- float4 FragColor0 [[color(0)]];
+ float3 m_location_0 [[color(0)]];
};
fragment main0_out main0()
float2 FragColor1 = {};
FragColor0 = 1.0;
FragColor1 = float2(2.0, 3.0);
- out.FragColor0.x = FragColor0;
- out.FragColor0.yz = FragColor1;
+ out.m_location_0.x = FragColor0;
+ out.m_location_0.yz = FragColor1;
return out;
}
struct main0_in
{
- float4 Foo3 [[attribute(0)]];
+ float4 m_location_0 [[attribute(0)]];
};
vertex main0_out main0(main0_in in [[stage_in]])
main0_out out = {};
float3 Foo3 = {};
float Foo1 = {};
- Foo3 = in.Foo3.xyz;
- Foo1 = in.Foo3.w;
+ Foo3 = in.m_location_0.xyz;
+ Foo1 = in.m_location_0.w;
out.gl_Position = float4(Foo3, Foo1);
out.Foo = Foo3 + float3(Foo1);
return out;
--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float4 m_location_0 [[color(0)]];
+ float4 m_location_1 [[color(1)]];
+ float4 m_location_2 [[color(2)]];
+};
+
+struct main0_in
+{
+ float InC_0 [[user(locn0_1), flat]];
+ float InA_0 [[user(locn1), flat]];
+ float InC_1 [[user(locn1_1), flat]];
+ float2 InB_0 [[user(locn1_2), flat]];
+ float InA_1 [[user(locn2), flat]];
+ float InC_2 [[user(locn2_1), flat]];
+ float2 InB_1 [[user(locn2_2), flat]];
+ float InD [[user(locn3_1), sample_perspective]];
+ float InE [[user(locn4_2), center_no_perspective]];
+ float InF [[user(locn5_3), centroid_perspective]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+ main0_out out = {};
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ float D = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ InA[0] = in.InA_0;
+ InA[1] = in.InA_1;
+ InB[0] = in.InB_0;
+ InB[1] = in.InB_1;
+ InC[0] = in.InC_0;
+ InC[1] = in.InC_1;
+ InC[2] = in.InC_2;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = (in.InD + in.InE) + in.InF;
+ out.m_location_1.x = A[0];
+ out.m_location_2.x = A[1];
+ out.m_location_1.zw = B[0];
+ out.m_location_2.zw = B[1];
+ out.m_location_0.y = C[0];
+ out.m_location_1.y = C[1];
+ out.m_location_2.y = C[2];
+ out.m_location_0.w = D;
+ return out;
+}
+
--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float4 m_location_0;
+ float4 m_location_1;
+ float4 m_location_2;
+ float4 gl_Position;
+};
+
+struct main0_in
+{
+ float4 m_location_0 [[attribute(0)]];
+ float4 m_location_1 [[attribute(1)]];
+ float4 m_location_2 [[attribute(2)]];
+ float4 Pos [[attribute(4)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
+{
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ float D = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ float InD = {};
+ device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
+ InA[0] = in.m_location_1.x;
+ InA[1] = in.m_location_2.x;
+ InB[0] = in.m_location_1.zw;
+ InB[1] = in.m_location_2.zw;
+ InC[0] = in.m_location_0.y;
+ InC[1] = in.m_location_1.y;
+ InC[2] = in.m_location_2.y;
+ InD = in.m_location_0.w;
+ if (any(gl_GlobalInvocationID >= spvStageInputSize))
+ return;
+ out.gl_Position = in.Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = InD;
+ out.m_location_1.x = A[0];
+ out.m_location_2.x = A[1];
+ out.m_location_1.zw = B[0];
+ out.m_location_2.zw = B[1];
+ out.m_location_0.y = C[0];
+ out.m_location_1.y = C[1];
+ out.m_location_2.y = C[2];
+ out.m_location_0.w = D;
+}
+
--- /dev/null
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+ T elements[Num ? Num : 1];
+
+ thread T& operator [] (size_t pos) thread
+ {
+ return elements[pos];
+ }
+ constexpr const thread T& operator [] (size_t pos) const thread
+ {
+ return elements[pos];
+ }
+
+ device T& operator [] (size_t pos) device
+ {
+ return elements[pos];
+ }
+ constexpr const device T& operator [] (size_t pos) const device
+ {
+ return elements[pos];
+ }
+
+ constexpr const constant T& operator [] (size_t pos) const constant
+ {
+ return elements[pos];
+ }
+
+ threadgroup T& operator [] (size_t pos) threadgroup
+ {
+ return elements[pos];
+ }
+ constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+ {
+ return elements[pos];
+ }
+};
+
+struct main0_out
+{
+ float C_0 [[user(locn0_1)]];
+ float D [[user(locn0_3)]];
+ float A_0 [[user(locn1)]];
+ float C_1 [[user(locn1_1)]];
+ float2 B_0 [[user(locn1_2)]];
+ float A_1 [[user(locn2)]];
+ float C_2 [[user(locn2_1)]];
+ float2 B_1 [[user(locn2_2)]];
+ float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+ float4 m_location_0 [[attribute(0)]];
+ float4 m_location_1 [[attribute(1)]];
+ float4 m_location_2 [[attribute(2)]];
+ float4 Pos [[attribute(4)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+ main0_out out = {};
+ spvUnsafeArray<float, 2> A = {};
+ spvUnsafeArray<float2, 2> B = {};
+ spvUnsafeArray<float, 3> C = {};
+ spvUnsafeArray<float, 2> InA = {};
+ spvUnsafeArray<float2, 2> InB = {};
+ spvUnsafeArray<float, 3> InC = {};
+ float InD = {};
+ InA[0] = in.m_location_1.x;
+ InA[1] = in.m_location_2.x;
+ InB[0] = in.m_location_1.zw;
+ InB[1] = in.m_location_2.zw;
+ InC[0] = in.m_location_0.y;
+ InC[1] = in.m_location_1.y;
+ InC[2] = in.m_location_2.y;
+ InD = in.m_location_0.w;
+ out.gl_Position = in.Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ out.D = InD;
+ out.A_0 = A[0];
+ out.A_1 = A[1];
+ out.B_0 = B[0];
+ out.B_1 = B[1];
+ out.C_0 = C[0];
+ out.C_1 = C[1];
+ out.C_2 = C[2];
+ return out;
+}
+
--- /dev/null
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) flat in float InA[2];
+layout(location = 1, component = 2) flat in vec2 InB[2];
+layout(location = 0, component = 1) flat in float InC[3];
+layout(location = 3, component = 1) sample in float InD;
+layout(location = 4, component = 2) noperspective in float InE;
+layout(location = 5, component = 3) centroid in float InF;
+
+void main()
+{
+ A = InA;
+ B = InB;
+ C = InC;
+ D = InD + InE + InF;
+}
--- /dev/null
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) in float InA[2];
+layout(location = 1, component = 2) in vec2 InB[2];
+layout(location = 0, component = 1) in float InC[3];
+layout(location = 0, component = 3) in float InD;
+layout(location = 4) in vec4 Pos;
+
+void main()
+{
+ gl_Position = Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = InD;
+}
--- /dev/null
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) in float InA[2];
+layout(location = 1, component = 2) in vec2 InB[2];
+layout(location = 0, component = 1) in float InC[3];
+layout(location = 0, component = 3) in float InD;
+layout(location = 4) in vec4 Pos;
+
+void main()
+{
+ gl_Position = Pos;
+ A = InA;
+ B = InB;
+ C = InC;
+ D = InD;
+}
return new_type_id;
}
+bool CompilerMSL::add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+ SPIRVariable &var,
+ const SPIRType &type,
+ InterfaceBlockMeta &meta)
+{
+ // Deal with Component decorations.
+ const InterfaceBlockMeta::LocationMeta *location_meta = nullptr;
+ uint32_t location = ~0u;
+ if (has_decoration(var.self, DecorationLocation))
+ {
+ location = get_decoration(var.self, DecorationLocation);
+ auto location_meta_itr = meta.location_meta.find(location);
+ if (location_meta_itr != end(meta.location_meta))
+ location_meta = &location_meta_itr->second;
+ }
+
+ // Check if we need to pad fragment output to match a certain number of components.
+ if (location_meta)
+ {
+ bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
+ msl_options.pad_fragment_output_components &&
+ get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;
+
+ auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
+ uint32_t start_component = get_decoration(var.self, DecorationComponent);
+ uint32_t type_components = type.vecsize;
+ uint32_t num_components = location_meta->num_components;
+
+ if (pad_fragment_output)
+ {
+ uint32_t locn = get_decoration(var.self, DecorationLocation);
+ num_components = std::max(num_components, get_target_components_for_fragment_location(locn));
+ }
+
+ // We have already declared an IO block member as m_location_N.
+ // Just emit an early-declared variable and fixup as needed.
+ // Arrays need to be unrolled here since each location might need a different number of components.
+ entry_func.add_local_variable(var.self);
+ vars_needing_early_declaration.push_back(var.self);
+
+ if (var.storage == StorageClassInput)
+ {
+ entry_func.fixup_hooks_in.push_back([=, &type, &var]() {
+ if (!type.array.empty())
+ {
+ uint32_t array_size = to_array_size_literal(type);
+ for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
+ {
+ statement(to_name(var.self), "[", loc_off, "]", " = ", ib_var_ref,
+ ".m_location_", location + loc_off,
+ vector_swizzle(type_components, start_component), ";");
+ }
+ }
+ else
+ {
+ statement(to_name(var.self), " = ", ib_var_ref, ".m_location_", location,
+ vector_swizzle(type_components, start_component), ";");
+ }
+ });
+ }
+ else
+ {
+ entry_func.fixup_hooks_out.push_back([=, &type, &var]() {
+ if (!type.array.empty())
+ {
+ uint32_t array_size = to_array_size_literal(type);
+ for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
+ {
+ statement(ib_var_ref, ".m_location_", location + loc_off,
+ vector_swizzle(type_components, start_component), " = ",
+ to_name(var.self), "[", loc_off, "];");
+ }
+ }
+ else
+ {
+ statement(ib_var_ref, ".m_location_", location,
+ vector_swizzle(type_components, start_component), " = ", to_name(var.self), ";");
+ }
+ });
+ }
+ return true;
+ }
+ else
+ return false;
+}
+
void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta)
{
auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
- // Deal with Component decorations.
- InterfaceBlockMeta::LocationMeta *location_meta = nullptr;
- if (has_decoration(var.self, DecorationLocation))
- {
- auto location_meta_itr = meta.location_meta.find(get_decoration(var.self, DecorationLocation));
- if (location_meta_itr != end(meta.location_meta))
- location_meta = &location_meta_itr->second;
- }
+ if (add_component_variable_to_interface_block(storage, ib_var_ref, var, type, meta))
+ return;
bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
msl_options.pad_fragment_output_components &&
get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;
- // Check if we need to pad fragment output to match a certain number of components.
- if (location_meta)
- {
- start_component = get_decoration(var.self, DecorationComponent);
- uint32_t num_components = location_meta->num_components;
- if (pad_fragment_output)
- {
- uint32_t locn = get_decoration(var.self, DecorationLocation);
- num_components = std::max(num_components, get_target_components_for_fragment_location(locn));
- }
-
- if (location_meta->ib_index != ~0u)
- {
- // We have already declared the variable. Just emit an early-declared variable and fixup as needed.
- entry_func.add_local_variable(var.self);
- vars_needing_early_declaration.push_back(var.self);
-
- if (var.storage == StorageClassInput)
- {
- uint32_t ib_index = location_meta->ib_index;
- entry_func.fixup_hooks_in.push_back([=, &var]() {
- statement(to_name(var.self), " = ", ib_var_ref, ".", to_member_name(ib_type, ib_index),
- vector_swizzle(type_components, start_component), ";");
- });
- }
- else
- {
- uint32_t ib_index = location_meta->ib_index;
- entry_func.fixup_hooks_out.push_back([=, &var]() {
- statement(ib_var_ref, ".", to_member_name(ib_type, ib_index),
- vector_swizzle(type_components, start_component), " = ", to_name(var.self), ";");
- });
- }
- return;
- }
- else
- {
- location_meta->ib_index = uint32_t(ib_type.member_types.size());
- type_id = build_extended_vector_type(type_id, num_components);
- if (var.storage == StorageClassInput)
- padded_input = true;
- else
- padded_output = true;
- }
- }
- else if (pad_fragment_output)
+ if (pad_fragment_output)
{
uint32_t locn = get_decoration(var.self, DecorationLocation);
target_components = get_target_components_for_fragment_location(locn);
uint32_t locn = get_decoration(var.self, DecorationLocation);
if (storage == StorageClassInput)
{
- type_id = ensure_correct_input_type(var.basetype, locn,
- location_meta ? location_meta->num_components : 0,
- meta.strip_array);
- if (!location_meta)
- var.basetype = type_id;
+ type_id = ensure_correct_input_type(var.basetype, locn, 0, meta.strip_array);
+ var.basetype = type_id;
type_id = get_pointee_type_id(type_id);
if (meta.strip_array && is_array(get<SPIRType>(type_id)))
mark_location_as_used_by_shader(locn, type, storage);
}
- if (!location_meta)
+ if (get_decoration_bitset(var.self).get(DecorationComponent))
{
- if (get_decoration_bitset(var.self).get(DecorationComponent))
- {
- uint32_t component = get_decoration(var.self, DecorationComponent);
- set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, component);
- }
+ uint32_t component = get_decoration(var.self, DecorationComponent);
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, component);
}
if (get_decoration_bitset(var.self).get(DecorationIndex))
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
}
- // If we have location meta, there is no unique OrigID. We won't need it, since we flatten/unflatten
- // the variable to stack anyways here.
- if (!location_meta)
- set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
+ set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
}
void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var);
uint32_t elem_cnt = 0;
+ if (add_component_variable_to_interface_block(storage, ib_var_ref, var, var_type, meta))
+ return;
+
if (is_matrix(var_type))
{
if (is_array(var_type))
{
auto &location_meta = meta.location_meta[location + location_offset];
location_meta.num_components = std::max(location_meta.num_components, component + type.vecsize);
+
+ // For variables sharing location, decorations and base type must match.
+ location_meta.base_type_id = type.self;
+ location_meta.flat = has_decoration(var.self, DecorationFlat);
+ location_meta.noperspective = has_decoration(var.self, DecorationNoPerspective);
+ location_meta.centroid = has_decoration(var.self, DecorationCentroid);
+ location_meta.sample = has_decoration(var.self, DecorationSample);
}
}
}
}
}
+ // When multiple variables need to access same location,
+ // unroll locations one by one and we will flatten output or input as necessary.
+ for (auto &loc : meta.location_meta)
+ {
+ uint32_t location = loc.first;
+ auto &location_meta = loc.second;
+
+ uint32_t ib_mbr_idx = uint32_t(ib_type.member_types.size());
+ uint32_t type_id = build_extended_vector_type(location_meta.base_type_id, location_meta.num_components);
+ ib_type.member_types.push_back(type_id);
+
+ set_member_name(ib_type.self, ib_mbr_idx, join("m_location_", location));
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+ mark_location_as_used_by_shader(location, get<SPIRType>(type_id), storage);
+
+ if (location_meta.flat)
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationFlat);
+ if (location_meta.noperspective)
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationNoPerspective);
+ if (location_meta.centroid)
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationCentroid);
+ if (location_meta.sample)
+ set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
+ }
+
// Sort the members of the structure by their locations.
MemberSorter member_sorter(ib_type, ir.meta[ib_type_id], MemberSorter::LocationThenBuiltInType);
member_sorter.sort();
{
struct LocationMeta
{
+ uint32_t base_type_id = 0;
uint32_t num_components = 0;
- uint32_t ib_index = ~0u;
+ bool flat = false;
+ bool noperspective = false;
+ bool centroid = false;
+ bool sample = false;
};
std::unordered_map<uint32_t, LocationMeta> location_meta;
bool strip_array = false;
SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
void add_plain_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
+ bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+ SPIRVariable &var, const SPIRType &type,
+ InterfaceBlockMeta &meta);
void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
SPIRType &ib_type, SPIRVariable &var, uint32_t index,
InterfaceBlockMeta &meta);
projects / platform / upstream / SPIRV-Cross.git / commitdiff