Welcome to mirror list, hosted at ThFree Co, Russian Federation.

github.com/KhronosGroup/SPIRV-Cross.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorHans-Kristian Arntzen <hans-kristian.arntzen@arm.com>2018-11-01 16:56:25 +0300
committerHans-Kristian Arntzen <hans-kristian.arntzen@arm.com>2018-11-01 16:58:02 +0300
commit480acdad1868b41d6130f7e343b37bc7a58e28cf (patch)
tree4dfa4da710b6ec8d8217820c1b314b73633a1836
parent37dbdf14faf8ce62250500ac8b60f57400f20378 (diff)
Deal with OpSpecConstantOp used as array size.
When trying to validate buffer sizes, we usually need to bail out when using SpecConstantOps, but for some very specific cases where we allow unsized arrays currently, we can safely allow "unknown" sized arrays as well. This is probably the best we can do, when we have even more difficult cases than this, we throw a more sensible error message.
-rw-r--r--reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp49
-rw-r--r--reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp46
-rw-r--r--reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp46
-rw-r--r--reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk34
-rw-r--r--reference/shaders-hlsl/comp/spec-constant-op-member-array.comp49
-rw-r--r--reference/shaders-msl/comp/spec-constant-op-member-array.comp46
-rw-r--r--reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp46
-rw-r--r--reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk34
-rw-r--r--shaders-hlsl/comp/spec-constant-op-member-array.comp33
-rw-r--r--shaders-msl/comp/spec-constant-op-member-array.comp33
-rw-r--r--shaders/vulkan/comp/spec-constant-op-member-array.vk.comp33
-rw-r--r--spirv_glsl.cpp37
-rw-r--r--spirv_glsl.hpp1
-rw-r--r--spirv_msl.cpp40
-rw-r--r--spirv_msl.hpp1
15 files changed, 494 insertions, 34 deletions
diff --git a/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp b/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..ad815b28
--- /dev/null
+++ b/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,49 @@
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+static const int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+static const int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+static const int c = SPIRV_CROSS_CONSTANT_ID_2;
+static const int _18 = (c + 50);
+#ifndef SPIRV_CROSS_CONSTANT_ID_3
+#define SPIRV_CROSS_CONSTANT_ID_3 400
+#endif
+static const int e = SPIRV_CROSS_CONSTANT_ID_3;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+RWByteAddressBuffer _22 : register(u0);
+
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+ uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void comp_main()
+{
+ _22.Store(gl_GlobalInvocationID.x * 4 + 2800, uint(int(_22.Load(gl_GlobalInvocationID.x * 4 + 2800)) + (int(_22.Load(gl_GlobalInvocationID.x * 4 + 2400)) + e)));
+}
+
+[numthreads(1, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+ gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+ comp_main();
+}
diff --git a/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp b/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..35ebe159
--- /dev/null
+++ b/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,46 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+constant int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+constant int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+constant int c = SPIRV_CROSS_CONSTANT_ID_2;
+constant int _18 = (c + 50);
+constant int e_tmp [[function_constant(3)]];
+constant int e = is_function_constant_defined(e_tmp) ? e_tmp : 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+struct SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+};
+
+kernel void main0(device SSBO& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp b/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
new file mode 100644
index 00000000..9344a7a2
--- /dev/null
+++ b/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
@@ -0,0 +1,46 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+const int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+const int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+const int c = SPIRV_CROSS_CONSTANT_ID_2;
+const int _18 = (c + 50);
+#ifndef SPIRV_CROSS_CONSTANT_ID_3
+#define SPIRV_CROSS_CONSTANT_ID_3 400
+#endif
+const int e = SPIRV_CROSS_CONSTANT_ID_3;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(binding = 0, std430) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+} _22;
+
+void main()
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk b/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk
new file mode 100644
index 00000000..d9a2822f
--- /dev/null
+++ b/reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk
@@ -0,0 +1,34 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(constant_id = 0) const int a = 100;
+layout(constant_id = 1) const int b = 200;
+layout(constant_id = 2) const int c = 300;
+const int _18 = (c + 50);
+layout(constant_id = 3) const int e = 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(set = 1, binding = 0, std430) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+} _22;
+
+void main()
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp b/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..ad815b28
--- /dev/null
+++ b/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,49 @@
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+static const int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+static const int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+static const int c = SPIRV_CROSS_CONSTANT_ID_2;
+static const int _18 = (c + 50);
+#ifndef SPIRV_CROSS_CONSTANT_ID_3
+#define SPIRV_CROSS_CONSTANT_ID_3 400
+#endif
+static const int e = SPIRV_CROSS_CONSTANT_ID_3;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+RWByteAddressBuffer _22 : register(u0);
+
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+ uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void comp_main()
+{
+ _22.Store(gl_GlobalInvocationID.x * 4 + 2800, uint(int(_22.Load(gl_GlobalInvocationID.x * 4 + 2800)) + (int(_22.Load(gl_GlobalInvocationID.x * 4 + 2400)) + e)));
+}
+
+[numthreads(1, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+ gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+ comp_main();
+}
diff --git a/reference/shaders-msl/comp/spec-constant-op-member-array.comp b/reference/shaders-msl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..35ebe159
--- /dev/null
+++ b/reference/shaders-msl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,46 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+constant int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+constant int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+constant int c = SPIRV_CROSS_CONSTANT_ID_2;
+constant int _18 = (c + 50);
+constant int e_tmp [[function_constant(3)]];
+constant int e = is_function_constant_defined(e_tmp) ? e_tmp : 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+struct SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+};
+
+kernel void main0(device SSBO& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp b/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
new file mode 100644
index 00000000..9344a7a2
--- /dev/null
+++ b/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
@@ -0,0 +1,46 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+#ifndef SPIRV_CROSS_CONSTANT_ID_0
+#define SPIRV_CROSS_CONSTANT_ID_0 100
+#endif
+const int a = SPIRV_CROSS_CONSTANT_ID_0;
+#ifndef SPIRV_CROSS_CONSTANT_ID_1
+#define SPIRV_CROSS_CONSTANT_ID_1 200
+#endif
+const int b = SPIRV_CROSS_CONSTANT_ID_1;
+#ifndef SPIRV_CROSS_CONSTANT_ID_2
+#define SPIRV_CROSS_CONSTANT_ID_2 300
+#endif
+const int c = SPIRV_CROSS_CONSTANT_ID_2;
+const int _18 = (c + 50);
+#ifndef SPIRV_CROSS_CONSTANT_ID_3
+#define SPIRV_CROSS_CONSTANT_ID_3 400
+#endif
+const int e = SPIRV_CROSS_CONSTANT_ID_3;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(binding = 0, std430) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+} _22;
+
+void main()
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk b/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk
new file mode 100644
index 00000000..d9a2822f
--- /dev/null
+++ b/reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk
@@ -0,0 +1,34 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(constant_id = 0) const int a = 100;
+layout(constant_id = 1) const int b = 200;
+layout(constant_id = 2) const int c = 300;
+const int _18 = (c + 50);
+layout(constant_id = 3) const int e = 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(set = 1, binding = 0, std430) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[_18];
+} _22;
+
+void main()
+{
+ _22.w[gl_GlobalInvocationID.x] += (_22.v[gl_GlobalInvocationID.x] + e);
+}
+
diff --git a/shaders-hlsl/comp/spec-constant-op-member-array.comp b/shaders-hlsl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..0b428eb0
--- /dev/null
+++ b/shaders-hlsl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,33 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(constant_id = 0) const int a = 100;
+layout(constant_id = 1) const int b = 200;
+layout(constant_id = 2) const int c = 300;
+const int d = c + 50;
+layout(constant_id = 3) const int e = 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(set = 1, binding = 0) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[d];
+};
+
+void main()
+{
+ w[gl_GlobalInvocationID.x] += v[gl_GlobalInvocationID.x] + e;
+}
diff --git a/shaders-msl/comp/spec-constant-op-member-array.comp b/shaders-msl/comp/spec-constant-op-member-array.comp
new file mode 100644
index 00000000..0b428eb0
--- /dev/null
+++ b/shaders-msl/comp/spec-constant-op-member-array.comp
@@ -0,0 +1,33 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(constant_id = 0) const int a = 100;
+layout(constant_id = 1) const int b = 200;
+layout(constant_id = 2) const int c = 300;
+const int d = c + 50;
+layout(constant_id = 3) const int e = 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(set = 1, binding = 0) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[d];
+};
+
+void main()
+{
+ w[gl_GlobalInvocationID.x] += v[gl_GlobalInvocationID.x] + e;
+}
diff --git a/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp b/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
new file mode 100644
index 00000000..0b428eb0
--- /dev/null
+++ b/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp
@@ -0,0 +1,33 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(constant_id = 0) const int a = 100;
+layout(constant_id = 1) const int b = 200;
+layout(constant_id = 2) const int c = 300;
+const int d = c + 50;
+layout(constant_id = 3) const int e = 400;
+
+struct A
+{
+ int member0[a];
+ int member1[b];
+};
+
+struct B
+{
+ int member0[b];
+ int member1[a];
+};
+
+layout(set = 1, binding = 0) buffer SSBO
+{
+ A member_a;
+ B member_b;
+ int v[a];
+ int w[d];
+};
+
+void main()
+{
+ w[gl_GlobalInvocationID.x] += v[gl_GlobalInvocationID.x] + e;
+}
diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp
index 3b640afb..8370145d 100644
--- a/spirv_glsl.cpp
+++ b/spirv_glsl.cpp
@@ -1042,8 +1042,7 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f
{
if (!type.array.empty())
{
- return to_array_size_literal(type, uint32_t(type.array.size()) - 1) *
- type_to_packed_array_stride(type, flags, packing);
+ return to_array_size_literal(type) * type_to_packed_array_stride(type, flags, packing);
}
uint32_t size = 0;
@@ -1121,6 +1120,9 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin
uint32_t offset = 0;
uint32_t pad_alignment = 1;
+ bool is_top_level_block =
+ has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock);
+
for (uint32_t i = 0; i < type.member_types.size(); i++)
{
auto &memb_type = get<SPIRType>(type.member_types[i]);
@@ -1128,8 +1130,26 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin
// Verify alignment rules.
uint32_t packed_alignment = type_to_packed_alignment(memb_type, member_flags, packing);
- uint32_t packed_size = type_to_packed_size(memb_type, member_flags, packing);
+ // This is a rather dirty workaround to deal with some cases of OpSpecConstantOp used as array size, e.g:
+ // layout(constant_id = 0) const int s = 10;
+ // const int S = s + 5; // SpecConstantOp
+ // buffer Foo { int data[S]; }; // <-- Very hard for us to deduce a fixed value here,
+ // we would need full implementation of compile-time constant folding. :(
+ // If we are the last member of a struct, there might be cases where the actual size of that member is irrelevant
+ // for our analysis (e.g. unsized arrays).
+ // This lets us simply ignore that there are spec constant op sized arrays in our buffers.
+ // Querying size of this member will fail, so just don't call it unless we have to.
+ //
+ // This is likely "best effort" we can support without going into unacceptably complicated workarounds.
+ bool member_can_be_unsized =
+ is_top_level_block && size_t(i + 1) == type.member_types.size() && !memb_type.array.empty();
+
+ uint32_t packed_size = 0;
+ if (!member_can_be_unsized)
+ packed_size = type_to_packed_size(memb_type, member_flags, packing);
+
+ // We only need to care about this if we have non-array types which can straddle the vec4 boundary.
if (packing_is_hlsl(packing))
{
// If a member straddles across a vec4 boundary, alignment is actually vec4.
@@ -8574,6 +8594,11 @@ string CompilerGLSL::pls_decl(const PlsRemap &var)
to_name(variable.self));
}
+uint32_t CompilerGLSL::to_array_size_literal(const SPIRType &type) const
+{
+ return to_array_size_literal(type, uint32_t(type.array.size() - 1));
+}
+
uint32_t CompilerGLSL::to_array_size_literal(const SPIRType &type, uint32_t index) const
{
assert(type.array.size() == type.array_size_literal.size());
@@ -8587,6 +8612,12 @@ uint32_t CompilerGLSL::to_array_size_literal(const SPIRType &type, uint32_t inde
// Use the default spec constant value.
// This is the best we can do.
uint32_t array_size_id = type.array[index];
+
+ // Explicitly check for this case. The error message you would get (bad cast) makes no sense otherwise.
+ if (ir.ids[array_size_id].get_type() == TypeConstantOp)
+ SPIRV_CROSS_THROW("An array size was found to be an OpSpecConstantOp. This is not supported since "
+ "SPIRV-Cross cannot deduce the actual size here.");
+
uint32_t array_size = get<SPIRConstant>(array_size_id).scalar();
return array_size;
}
diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp
index 54e8bdb4..71dbc434 100644
--- a/spirv_glsl.hpp
+++ b/spirv_glsl.hpp
@@ -326,6 +326,7 @@ protected:
std::string type_to_array_glsl(const SPIRType &type);
std::string to_array_size(const SPIRType &type, uint32_t index);
uint32_t to_array_size_literal(const SPIRType &type, uint32_t index) const;
+ uint32_t to_array_size_literal(const SPIRType &type) const;
std::string variable_decl(const SPIRVariable &variable);
std::string variable_decl_function_local(SPIRVariable &variable);
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index aabb493b..63b95b7c 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -534,21 +534,6 @@ void CompilerMSL::localize_global_variables()
}
}
-// Metal does not allow dynamic array lengths.
-// Turn off specialization of any constants that are used for array lengths.
-void CompilerMSL::resolve_specialized_array_lengths()
-{
- for (auto &id : ir.ids)
- {
- if (id.get_type() == TypeConstant)
- {
- auto &c = id.get<SPIRConstant>();
- if (c.is_used_as_array_length)
- c.specialization = false;
- }
- }
-}
-
// For any global variable accessed directly by a function,
// extract that variable and add it as an argument to that function.
void CompilerMSL::extract_global_variables_from_functions()
@@ -1024,8 +1009,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage)
if (type.array.size() != 1)
SPIRV_CROSS_THROW("MSL cannot emit arrays-of-arrays in input and output variables.");
- elem_cnt = type.array_size_literal.back() ? type.array.back() :
- get<SPIRConstant>(type.array.back()).scalar();
+ elem_cnt = to_array_size_literal(type);
}
auto *usable_type = &type;
@@ -1199,14 +1183,13 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
MemberSorter member_sorter(ib_type, ir.meta[ib_type_id], MemberSorter::Offset);
member_sorter.sort();
- uint32_t curr_offset;
uint32_t mbr_cnt = uint32_t(ib_type.member_types.size());
// Test the alignment of each member, and if a member should be closer to the previous
// member than the default spacing expects, it is likely that the previous member is in
// a packed format. If so, and the previous member is packable, pack it.
// For example...this applies to any 3-element vector that is followed by a scalar.
- curr_offset = 0;
+ uint32_t curr_offset = 0;
for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++)
{
if (is_member_packable(ib_type, mbr_idx))
@@ -1228,7 +1211,9 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
}
// Increment the current offset to be positioned immediately after the current member.
- curr_offset = mbr_offset + uint32_t(get_declared_struct_member_size(ib_type, mbr_idx));
+ // Don't do this for the last member since it can be unsized, and it is not relevant for padding purposes here.
+ if (mbr_idx + 1 < mbr_cnt)
+ curr_offset = mbr_offset + uint32_t(get_declared_struct_member_size(ib_type, mbr_idx));
}
}
@@ -1259,7 +1244,7 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
uint32_t md_elem_cnt = 1;
size_t last_elem_idx = mbr_type.array.size() - 1;
for (uint32_t i = 0; i < last_elem_idx; i++)
- md_elem_cnt *= max(to_array_size_literal(mbr_type, i), 1U);
+ md_elem_cnt *= max(to_array_size_literal(mbr_type, i), 1u);
uint32_t unpacked_array_stride = unpacked_mbr_size * md_elem_cnt;
uint32_t array_stride = type_struct_member_array_stride(ib_type, index);
@@ -3909,8 +3894,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
// Metal doesn't directly support this, so we must expand the
// array. We'll declare a local array to hold these elements
// later.
- uint32_t array_size =
- type.array_size_literal.back() ? type.array.back() : get<SPIRConstant>(type.array.back()).scalar();
+ uint32_t array_size = to_array_size_literal(type);
if (array_size == 0)
SPIRV_CROSS_THROW("Unsized arrays of buffers are not supported in MSL.");
@@ -4332,8 +4316,7 @@ std::string CompilerMSL::sampler_type(const SPIRType &type)
SPIRV_CROSS_THROW("Arrays of arrays of samplers are not supported in MSL.");
// Arrays of samplers in MSL must be declared with a special array<T, N> syntax ala C++11 std::array.
- uint32_t array_size =
- type.array_size_literal.back() ? type.array.back() : get<SPIRConstant>(type.array.back()).scalar();
+ uint32_t array_size = to_array_size_literal(type);
if (array_size == 0)
SPIRV_CROSS_THROW("Unsized array of samplers is not supported in MSL.");
@@ -4375,8 +4358,7 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id)
SPIRV_CROSS_THROW("Arrays of arrays of textures are not supported in MSL.");
// Arrays of images in MSL must be declared with a special array<T, N> syntax ala C++11 std::array.
- uint32_t array_size =
- type.array_size_literal.back() ? type.array.back() : get<SPIRConstant>(type.array.back()).scalar();
+ uint32_t array_size = to_array_size_literal(type);
if (array_size == 0)
SPIRV_CROSS_THROW("Unsized array of images is not supported in MSL.");
@@ -4757,9 +4739,7 @@ size_t CompilerMSL::get_declared_struct_member_size(const SPIRType &struct_type,
// Runtime arrays will have zero size so force to min of one.
if (!type.array.empty())
{
- bool array_size_literal = type.array_size_literal.back();
- uint32_t array_size =
- array_size_literal ? type.array.back() : get<SPIRConstant>(type.array.back()).scalar();
+ uint32_t array_size = to_array_size_literal(type);
return type_struct_member_array_stride(struct_type, index) * max(array_size, 1u);
}
diff --git a/spirv_msl.hpp b/spirv_msl.hpp
index cc83b715..83027edf 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -341,7 +341,6 @@ protected:
void preprocess_op_codes();
void localize_global_variables();
void extract_global_variables_from_functions();
- void resolve_specialized_array_lengths();
void mark_packable_structs();
void mark_as_packable(SPIRType &type);