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

Deal with OpSpecConstantOp used as array size.

Browse Source 
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.
This commit is contained in:
Hans-Kristian Arntzen 2018-11-01 14:56:25 +01:00
parent 37dbdf14fa
commit 480acdad18
15 changed files with 494 additions and 34 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

49
reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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();
}

46
reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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);
}

46
reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp Normal file
View File

@ -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);
}

34
reference/opt/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk Normal file
View File

@ -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);
}

49
reference/shaders-hlsl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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();
}

46
reference/shaders-msl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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);
}

46
reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp Normal file
View File

@ -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);
}

34
reference/shaders/vulkan/comp/spec-constant-op-member-array.vk.comp.vk Normal file
View File

@ -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);
}

33
shaders-hlsl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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;
}

33
shaders-msl/comp/spec-constant-op-member-array.comp Normal file
View File

@ -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;
}

33
shaders/vulkan/comp/spec-constant-op-member-array.vk.comp Normal file
View File

@ -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;
}

37
spirv_glsl.cpp
View File

@ -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;
}

1
spirv_glsl.hpp
View File

@ -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);

40
spirv_msl.cpp
View File

@ -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);
}

1
spirv_msl.hpp
View File

@ -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);