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

Merge pull request #1376 from KhronosGroup/fix-1374

Browse Source 
Handle physical pointers in reflection API.
This commit is contained in:
Hans-Kristian Arntzen 2020-05-25 15:07:22 +02:00 committed by GitHub
commit 61cddd6307
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
13 changed files with 393 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

71
reference/shaders-reflection/asm/comp/pointer-to-array-of-physical-pointer.asm.comp.json Normal file
View File

@ -0,0 +1,71 @@
{
"entryPoints" : [
{
"name" : "main",
"mode" : "comp",
"workgroup_size" : [
1,
1,
1
],
"workgroup_size_is_spec_constant_id" : [
false,
false,
false
]
}
],
"types" : {
"_3" : {
"name" : "Params",
"members" : [
{
"name" : "x",
"type" : "float",
"offset" : 0
},
{
"name" : "y",
"type" : "_6",
"offset" : 16,
"physical_pointer" : true
}
]
},
"_4" : {
"name" : "IntBuf",
"members" : [
{
"name" : "v",
"type" : "int",
"offset" : 0
}
]
},
"_11" : {
"name" : "IntBuf",
"type" : "_4",
"physical_pointer" : true
},
"_6" : {
"name" : "IntBuf",
"array" : [
3
],
"array_size_is_literal" : [
true
],
"type" : "_11",
"array_stride" : 16
}
},
"ubos" : [
{
"type" : "_3",
"name" : "Params",
"block_size" : 24,
"set" : 0,
"binding" : 0
}
]
}

66
reference/shaders-reflection/comp/array-of-physical-pointer.comp.json Normal file
View File

@ -0,0 +1,66 @@
{
"entryPoints" : [
{
"name" : "main",
"mode" : "comp",
"workgroup_size" : [
1,
1,
1
],
"workgroup_size_is_spec_constant_id" : [
false,
false,
false
]
}
],
"types" : {
"_11" : {
"name" : "Params",
"members" : [
{
"name" : "x",
"type" : "float",
"offset" : 0
},
{
"name" : "y",
"type" : "_7",
"array" : [
3
],
"array_size_is_literal" : [
true
],
"offset" : 16,
"array_stride" : 16
}
]
},
"_13" : {
"name" : "IntBuf",
"members" : [
{
"name" : "v",
"type" : "int",
"offset" : 0
}
]
},
"_7" : {
"name" : "IntBuf",
"type" : "_13",
"physical_pointer" : true
}
},
"ubos" : [
{
"type" : "_11",
"name" : "Params",
"block_size" : 64,
"set" : 0,
"binding" : 0
}
]
}

55
reference/shaders-reflection/comp/physical-pointer.comp.json Normal file
View File

@ -0,0 +1,55 @@
{
"entryPoints" : [
{
"name" : "main",
"mode" : "comp",
"workgroup_size" : [
1,
1,
1
],
"workgroup_size_is_spec_constant_id" : [
false,
false,
false
]
}
],
"types" : {
"_8" : {
"name" : "Params",
"members" : [
{
"name" : "x",
"type" : "float",
"offset" : 0
},
{
"name" : "y",
"type" : "_10",
"offset" : 8,
"physical_pointer" : true
}
]
},
"_10" : {
"name" : "IntBuf",
"members" : [
{
"name" : "v",
"type" : "int",
"offset" : 0
}
]
}
},
"ubos" : [
{
"type" : "_8",
"name" : "Params",
"block_size" : 16,
"set" : 0,
"binding" : 0
}
]
}

51
shaders-reflection/asm/comp/pointer-to-array-of-physical-pointer.asm.comp Normal file
View File

@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 8
; Bound: 17
; Schema: 0
OpCapability Shader
OpCapability PhysicalStorageBufferAddresses
OpExtension "SPV_EXT_physical_storage_buffer"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel PhysicalStorageBuffer64 GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 460
OpSourceExtension "GL_EXT_buffer_reference"
OpSourceExtension "GL_EXT_buffer_reference2"
OpSourceExtension "GL_EXT_shader_explicit_arithmetic_types_int64"
OpName %main "main"
OpName %Params "Params"
OpMemberName %Params 0 "x"
OpMemberName %Params 1 "y"
OpName %IntBuf "IntBuf"
OpMemberName %IntBuf 0 "v"
OpName %_ ""
OpDecorate %_arr_7_uint_3 ArrayStride 16
OpMemberDecorate %Params 0 Offset 0
OpMemberDecorate %Params 1 Offset 16
OpDecorate %Params Block
OpMemberDecorate %IntBuf 0 Offset 0
OpDecorate %IntBuf Block
OpDecorate %_arr__ptr_PhysicalStorageBuffer_IntBuf_uint_3 ArrayStride 16
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
OpTypeForwardPointer %_ptr_PhysicalStorageBuffer_IntBuf PhysicalStorageBuffer
%uint = OpTypeInt 32 0
%uint_3 = OpConstant %uint 3
%_arr_7_uint_3 = OpTypeArray %_ptr_PhysicalStorageBuffer_IntBuf %uint_3
%ptr_array_ptr = OpTypePointer PhysicalStorageBuffer %_arr_7_uint_3
%Params = OpTypeStruct %float %ptr_array_ptr
%int = OpTypeInt 32 1
%IntBuf = OpTypeStruct %int
%_ptr_PhysicalStorageBuffer_IntBuf = OpTypePointer PhysicalStorageBuffer %IntBuf
%_arr__ptr_PhysicalStorageBuffer_IntBuf_uint_3 = OpTypeArray %_ptr_PhysicalStorageBuffer_IntBuf %uint_3
%_ptr_Uniform_Params = OpTypePointer Uniform %Params
%_ = OpVariable %_ptr_Uniform_Params Uniform
%main = OpFunction %void None %3
%5 = OpLabel
OpReturn
OpFunctionEnd

15
shaders-reflection/comp/array-of-physical-pointer.comp Normal file
View File

@ -0,0 +1,15 @@
#version 460
#extension GL_EXT_shader_explicit_arithmetic_types_int64 : enable
#extension GL_EXT_buffer_reference2 : enable
layout(buffer_reference, std430, buffer_reference_align = 4) buffer IntBuf
{
int v;
};
layout(std140, binding = 0) uniform Params
{
float x;
IntBuf y[3];
};
void main()
{
}

15
shaders-reflection/comp/physical-pointer.comp Normal file
View File

@ -0,0 +1,15 @@
#version 460
#extension GL_EXT_shader_explicit_arithmetic_types_int64 : enable
#extension GL_EXT_buffer_reference2 : enable
layout(buffer_reference, std430, buffer_reference_align = 4) buffer IntBuf
{
int v;
};
layout(std140, binding = 0) uniform Params
{
float x;
IntBuf y;
};
void main()
{
}

1
spirv_common.hpp
View File

@ -558,6 +558,7 @@ struct SPIRType : IVariant
// Keep track of how many pointer layers we have.
uint32_t pointer_depth = 0;
bool pointer = false;
bool forward_pointer = false;
spv::StorageClass storage = spv::StorageClassGeneric;

13
spirv_cross.cpp
View File

@ -1659,6 +1659,13 @@ size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, ui
break;
}
if (type.pointer && type.storage == StorageClassPhysicalStorageBuffer)
{
// Check if this is a top-level pointer type, and not an array of pointers.
if (type.pointer_depth > get<SPIRType>(type.parent_type).pointer_depth)
return 8;
}
if (!type.array.empty())
{
// For arrays, we can use ArrayStride to get an easy check.
@ -4637,6 +4644,12 @@ bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
return type.pointer_depth == get<SPIRType>(type.parent_type).pointer_depth;
}
bool Compiler::type_is_top_level_physical_pointer(const SPIRType &type) const
{
return type.pointer && type.storage == StorageClassPhysicalStorageBuffer &&
type.pointer_depth > get<SPIRType>(type.parent_type).pointer_depth;
}
bool Compiler::flush_phi_required(BlockID from, BlockID to) const
{
auto &child = get<SPIRBlock>(to);

1
spirv_cross.hpp
View File

@ -1037,6 +1037,7 @@ protected:
void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration);
bool type_is_array_of_pointers(const SPIRType &type) const;
bool type_is_top_level_physical_pointer(const SPIRType &type) const;
bool type_is_block_like(const SPIRType &type) const;
bool type_is_opaque_value(const SPIRType &type) const;

24
spirv_parser.cpp
View File

@ -119,6 +119,16 @@ void Parser::parse()
for (auto &i : instructions)
parse(i);
for (auto &fixup : forward_pointer_fixups)
{
auto &target = get<SPIRType>(fixup.first);
auto &source = get<SPIRType>(fixup.second);
target.member_types = source.member_types;
target.basetype = source.basetype;
target.self = source.self;
}
forward_pointer_fixups.clear();
if (current_function)
SPIRV_CROSS_THROW("Function was not terminated.");
if (current_block)
@ -543,6 +553,11 @@ void Parser::parse(const Instruction &instruction)
auto *c = maybe_get<SPIRConstant>(cid);
bool literal = c && !c->specialization;
// We're copying type information into Array types, so we'll need a fixup for any physical pointer
// references.
if (base.forward_pointer)
forward_pointer_fixups.push_back({ id, tid });
arraybase.array_size_literal.push_back(literal);
arraybase.array.push_back(literal ? c->scalar() : cid);
// Do NOT set arraybase.self!
@ -556,6 +571,11 @@ void Parser::parse(const Instruction &instruction)
auto &base = get<SPIRType>(ops[1]);
auto &arraybase = set<SPIRType>(id);
// We're copying type information into Array types, so we'll need a fixup for any physical pointer
// references.
if (base.forward_pointer)
forward_pointer_fixups.push_back({ id, ops[1] });
arraybase = base;
arraybase.array.push_back(0);
arraybase.array_size_literal.push_back(true);
@ -614,6 +634,9 @@ void Parser::parse(const Instruction &instruction)
if (ptrbase.storage == StorageClassAtomicCounter)
ptrbase.basetype = SPIRType::AtomicCounter;
if (base.forward_pointer)
forward_pointer_fixups.push_back({ id, ops[2] });
ptrbase.parent_type = ops[2];
// Do NOT set ptrbase.self!
@ -627,6 +650,7 @@ void Parser::parse(const Instruction &instruction)
ptrbase.pointer = true;
ptrbase.pointer_depth++;
ptrbase.storage = static_cast<StorageClass>(ops[1]);
ptrbase.forward_pointer = true;
if (ptrbase.storage == StorageClassAtomicCounter)
ptrbase.basetype = SPIRType::AtomicCounter;

1
spirv_parser.hpp
View File

@ -84,6 +84,7 @@ private:
// This must be an ordered data structure so we always pick the same type aliases.
SmallVector<uint32_t> global_struct_cache;
SmallVector<std::pair<uint32_t, uint32_t>> forward_pointer_fixups;
bool types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const;
bool variable_storage_is_aliased(const SPIRVariable &v) const;

101
spirv_reflect.cpp
View File

@ -277,23 +277,54 @@ string CompilerReflection::compile()
return json_stream->str();
}
static bool naturally_emit_type(const SPIRType &type)
{
return type.basetype == SPIRType::Struct && !type.pointer && type.array.empty();
}
bool CompilerReflection::type_is_reference(const SPIRType &type) const
{
// Physical pointers and arrays of physical pointers need to refer to the pointee's type.
return type_is_top_level_physical_pointer(type) ||
(!type.array.empty() && type_is_top_level_physical_pointer(get<SPIRType>(type.parent_type)));
}
void CompilerReflection::emit_types()
{
bool emitted_open_tag = false;
ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
if (type.basetype == SPIRType::Struct && !type.pointer && type.array.empty())
emit_type(type, emitted_open_tag);
SmallVector<uint32_t> physical_pointee_types;
// If we have physical pointers or arrays of physical pointers, it's also helpful to emit the pointee type
// and chain the type hierarchy. For POD, arrays can emit the entire type in-place.
ir.for_each_typed_id<SPIRType>([&](uint32_t self, SPIRType &type) {
if (naturally_emit_type(type))
{
emit_type(self, emitted_open_tag);
}
else if (type_is_reference(type))
{
if (!naturally_emit_type(this->get<SPIRType>(type.parent_type)) &&
find(physical_pointee_types.begin(), physical_pointee_types.end(),
type.parent_type) == physical_pointee_types.end())
{
physical_pointee_types.push_back(type.parent_type);
}
}
});
for (uint32_t pointee_type : physical_pointee_types)
emit_type(pointee_type, emitted_open_tag);
if (emitted_open_tag)
{
json_stream->end_json_object();
}
}
void CompilerReflection::emit_type(const SPIRType &type, bool &emitted_open_tag)
void CompilerReflection::emit_type(uint32_t type_id, bool &emitted_open_tag)
{
auto &type = get<SPIRType>(type_id);
auto name = type_to_glsl(type);
if (type.type_alias != TypeID(0))
@ -304,26 +335,42 @@ void CompilerReflection::emit_type(const SPIRType &type, bool &emitted_open_tag)
json_stream->emit_json_key_object("types");
emitted_open_tag = true;
}
json_stream->emit_json_key_object("_" + std::to_string(type.self));
json_stream->emit_json_key_object("_" + std::to_string(type_id));
json_stream->emit_json_key_value("name", name);
json_stream->emit_json_key_array("members");
// FIXME ideally we'd like to emit the size of a structure as a
// convenience to people parsing the reflected JSON. The problem
// is that there's no implicit size for a type. It's final size
// will be determined by the top level declaration in which it's
// included. So there might be one size for the struct if it's
// included in a std140 uniform block and another if it's included
// in a std430 uniform block.
// The solution is to include *all* potential sizes as a map of
// layout type name to integer, but that will probably require
// some additional logic being written in this class, or in the
// parent CompilerGLSL class.
auto size = type.member_types.size();
for (uint32_t i = 0; i < size; ++i)
if (type_is_top_level_physical_pointer(type))
{
emit_type_member(type, i);
json_stream->emit_json_key_value("type", "_" + std::to_string(type.parent_type));
json_stream->emit_json_key_value("physical_pointer", true);
}
json_stream->end_json_array();
else if (!type.array.empty())
{
emit_type_array(type);
json_stream->emit_json_key_value("type", "_" + std::to_string(type.parent_type));
json_stream->emit_json_key_value("array_stride", get_decoration(type_id, DecorationArrayStride));
}
else
{
json_stream->emit_json_key_array("members");
// FIXME ideally we'd like to emit the size of a structure as a
// convenience to people parsing the reflected JSON. The problem
// is that there's no implicit size for a type. It's final size
// will be determined by the top level declaration in which it's
// included. So there might be one size for the struct if it's
// included in a std140 uniform block and another if it's included
// in a std430 uniform block.
// The solution is to include *all* potential sizes as a map of
// layout type name to integer, but that will probably require
// some additional logic being written in this class, or in the
// parent CompilerGLSL class.
auto size = type.member_types.size();
for (uint32_t i = 0; i < size; ++i)
{
emit_type_member(type, i);
}
json_stream->end_json_array();
}
json_stream->end_json_object();
}
@ -335,7 +382,12 @@ void CompilerReflection::emit_type_member(const SPIRType &type, uint32_t index)
// FIXME we'd like to emit the offset of each member, but such offsets are
// context dependent. See the comment above regarding structure sizes
json_stream->emit_json_key_value("name", name);
if (membertype.basetype == SPIRType::Struct)
if (type_is_reference(membertype))
{
json_stream->emit_json_key_value("type", "_" + std::to_string(membertype.parent_type));
}
else if (membertype.basetype == SPIRType::Struct)
{
json_stream->emit_json_key_value("type", "_" + std::to_string(membertype.self));
}
@ -349,7 +401,7 @@ void CompilerReflection::emit_type_member(const SPIRType &type, uint32_t index)
void CompilerReflection::emit_type_array(const SPIRType &type)
{
if (!type.array.empty())
if (!type_is_top_level_physical_pointer(type) && !type.array.empty())
{
json_stream->emit_json_key_array("array");
// Note that we emit the zeros here as a means of identifying
@ -388,6 +440,9 @@ void CompilerReflection::emit_type_member_qualifiers(const SPIRType &type, uint3
json_stream->emit_json_key_value("matrix_stride", dec.matrix_stride);
if (dec.decoration_flags.get(DecorationRowMajor))
json_stream->emit_json_key_value("row_major", true);
if (type_is_top_level_physical_pointer(membertype))
json_stream->emit_json_key_value("physical_pointer", true);
}
}

3
spirv_reflect.hpp
View File

@ -67,11 +67,12 @@ private:
void emit_resources();
void emit_specialization_constants();
void emit_type(const SPIRType &type, bool &emitted_open_tag);
void emit_type(uint32_t type_id, bool &emitted_open_tag);
void emit_type_member(const SPIRType &type, uint32_t index);
void emit_type_member_qualifiers(const SPIRType &type, uint32_t index);
void emit_type_array(const SPIRType &type);
void emit_resources(const char *tag, const SmallVector<Resource> &resources);
bool type_is_reference(const SPIRType &type) const;
std::string to_member_name(const SPIRType &type, uint32_t index) const;