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

Merge pull request #737 from KhronosGroup/fix-724

Browse Source 
Deal with OpSpecConstantOp used as array size.
This commit is contained in:
Hans-Kristian Arntzen 2018-11-01 15:26:19 +01:00 committed by GitHub
commit 13633c0500
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
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);