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

Add basic support for StorageClassStorageBuffer.

Browse Source 
Needs more testing, but seems to work.
This commit is contained in:
Hans-Kristian Arntzen 2017-09-28 13:28:44 +02:00
parent fae64f032d
commit 153fed031d
6 changed files with 70 additions and 44 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
spirv_common.hpp
View File

@ -713,14 +713,16 @@ struct SPIRConstant : IVariant
struct ConstantVector
{
Constant r[4];
uint32_t id[4] = {}; // If != 0, this element is a specialization constant, and we should keep track of it as such.
// If != 0, this element is a specialization constant, and we should keep track of it as such.
uint32_t id[4] = {};
uint32_t vecsize = 1;
};
struct ConstantMatrix
{
ConstantVector c[4];
uint32_t id[4] = {}; // If != 0, this column is a specialization constant, and we should keep track of it as such.
// If != 0, this column is a specialization constant, and we should keep track of it as such.
uint32_t id[4] = {};
uint32_t columns = 1;
};
@ -793,7 +795,8 @@ struct SPIRConstant : IVariant
}
SPIRConstant(uint32_t constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized)
: constant_type(constant_type_), specialization(specialized)
: constant_type(constant_type_)
, specialization(specialized)
{
subconstants.insert(end(subconstants), elements, elements + num_elements);
specialization = specialized;
@ -801,7 +804,8 @@ struct SPIRConstant : IVariant
// Construct scalar (32-bit).
SPIRConstant(uint32_t constant_type_, uint32_t v0, bool specialized)
: constant_type(constant_type_), specialization(specialized)
: constant_type(constant_type_)
, specialization(specialized)
{
m.c[0].r[0].u32 = v0;
m.c[0].vecsize = 1;
@ -810,7 +814,8 @@ struct SPIRConstant : IVariant
// Construct scalar (64-bit).
SPIRConstant(uint32_t constant_type_, uint64_t v0, bool specialized)
: constant_type(constant_type_), specialization(specialized)
: constant_type(constant_type_)
, specialization(specialized)
{
m.c[0].r[0].u64 = v0;
m.c[0].vecsize = 1;
@ -818,8 +823,10 @@ struct SPIRConstant : IVariant
}
// Construct vectors and matrices.
SPIRConstant(uint32_t constant_type_, const SPIRConstant * const *vector_elements, uint32_t num_elements, bool specialized)
: constant_type(constant_type_), specialization(specialized)
SPIRConstant(uint32_t constant_type_, const SPIRConstant *const *vector_elements, uint32_t num_elements,
bool specialized)
: constant_type(constant_type_)
, specialization(specialized)
{
bool matrix = vector_elements[0]->m.c[0].vecsize > 1;

20
spirv_cross.cpp
View File

@ -86,7 +86,8 @@ string Compiler::compile()
bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
{
auto &type = get<SPIRType>(v.basetype);
bool ssbo = (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0;
bool ssbo = v.storage == StorageClassStorageBuffer ||
((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0);
bool image = type.basetype == SPIRType::Image;
bool counter = type.basetype == SPIRType::AtomicCounter;
bool is_restrict = (meta[v.self].decoration.decoration_flags & (1ull << DecorationRestrict)) != 0;
@ -426,6 +427,7 @@ static inline bool storage_class_is_interface(spv::StorageClass storage)
case StorageClassUniformConstant:
case StorageClassAtomicCounter:
case StorageClassPushConstant:
case StorageClassStorageBuffer:
return true;
default:
@ -655,12 +657,17 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *ac
{
res.uniform_buffers.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
}
// SSBOs
// Old way to declare SSBOs.
else if (type.storage == StorageClassUniform &&
(meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)))
{
res.storage_buffers.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
}
// Modern way to declare SSBOs.
else if (type.storage == StorageClassStorageBuffer)
{
res.storage_buffers.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
}
// Push constant blocks
else if (type.storage == StorageClassPushConstant)
{
@ -777,8 +784,7 @@ void Compiler::parse()
if (id.get_type() == TypeConstant)
{
auto &c = id.get<SPIRConstant>();
if (meta[c.self].decoration.builtin &&
meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
if (meta[c.self].decoration.builtin && meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
{
// In current SPIR-V, there can be just one constant like this.
// All entry points will receive the constant value.
@ -3549,7 +3555,8 @@ bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const
{
auto *var = maybe_get<SPIRVariable>(id);
// Ensure that this is actually a buffer object.
return var && has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock);
return var && (var->storage == StorageClassStorageBuffer ||
has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock));
}
else
return false;
@ -3565,7 +3572,8 @@ bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id)
{
auto *var = maybe_get<SPIRVariable>(i);
// Ensure that this is actually a buffer object.
if (var && has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock))
if (var && (var->storage == StorageClassStorageBuffer ||
has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock)))
{
counter_id = i;
return true;

5
spirv_cross.hpp
View File

@ -197,7 +197,7 @@ public:
// Returns a vector of which members of a struct are potentially in use by a
// SPIR-V shader. The granularity of this analysis is per-member of a struct.
// This can be used for Buffer (UBO), BufferBlock (SSBO) and PushConstant blocks.
// This can be used for Buffer (UBO), BufferBlock/StorageBuffer (SSBO) and PushConstant blocks.
// ID is the Resource::id obtained from get_shader_resources().
std::vector<BufferRange> get_active_buffer_ranges(uint32_t id) const;
@ -284,8 +284,7 @@ public:
// If the component is not a specialization constant, a zeroed out struct will be written.
// The return value is the constant ID of the builtin WorkGroupSize, but this is not expected to be useful
// for most use cases.
uint32_t get_work_group_size_specialization_constants(SpecializationConstant &x,
SpecializationConstant &y,
uint32_t get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
SpecializationConstant &z) const;
// Analyzes all separate image and samplers used from the currently selected entry point,

22
spirv_glsl.cpp
View File

@ -1056,7 +1056,8 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
// If SPIR-V does not comply with either layout, we cannot really work around it.
if (var.storage == StorageClassUniform && (typeflags & (1ull << DecorationBlock)))
attr.push_back("std140");
else if (var.storage == StorageClassUniform && (typeflags & (1ull << DecorationBufferBlock)))
else if (var.storage == StorageClassStorageBuffer ||
(var.storage == StorageClassUniform && (typeflags & (1ull << DecorationBufferBlock))))
attr.push_back(ssbo_is_std430_packing(type) ? "std430" : "std140");
else if (options.vulkan_semantics && var.storage == StorageClassPushConstant)
attr.push_back(ssbo_is_std430_packing(type) ? "std430" : "std140");
@ -1135,7 +1136,8 @@ void CompilerGLSL::emit_buffer_block(const SPIRVariable &var)
void CompilerGLSL::emit_buffer_block_legacy(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
bool ssbo = (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0;
bool ssbo = var.storage == StorageClassStorageBuffer ||
((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0);
if (ssbo)
SPIRV_CROSS_THROW("SSBOs not supported in legacy targets.");
@ -1155,7 +1157,8 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
auto &type = get<SPIRType>(var.basetype);
uint64_t flags = get_buffer_block_flags(var);
bool ssbo = (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0;
bool ssbo = var.storage == StorageClassStorageBuffer ||
((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0);
bool is_restrict = ssbo && (flags & (1ull << DecorationRestrict)) != 0;
bool is_writeonly = ssbo && (flags & (1ull << DecorationNonReadable)) != 0;
bool is_readonly = ssbo && (flags & (1ull << DecorationNonWritable)) != 0;
@ -1386,9 +1389,7 @@ void CompilerGLSL::emit_specialization_constant(const SPIRConstant &constant)
SpecializationConstant wg_x, wg_y, wg_z;
uint32_t workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
if (constant.self == workgroup_size_id ||
constant.self == wg_x.id ||
constant.self == wg_y.id ||
if (constant.self == workgroup_size_id || constant.self == wg_x.id || constant.self == wg_y.id ||
constant.self == wg_z.id)
{
// These specialization constants are implicitly declared by emitting layout() in;
@ -1751,9 +1752,12 @@ void CompilerGLSL::emit_resources()
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (var.storage != StorageClassFunction && type.pointer && type.storage == StorageClassUniform &&
!is_hidden_variable(var) && (meta[type.self].decoration.decoration_flags &
((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))))
bool is_block_storage = type.storage == StorageClassStorageBuffer || type.storage == StorageClassUniform;
bool has_block_flags = (meta[type.self].decoration.decoration_flags &
((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))) != 0;
if (var.storage != StorageClassFunction && type.pointer && is_block_storage && !is_hidden_variable(var) &&
has_block_flags)
{
emit_buffer_block(var);
}

24
spirv_hlsl.cpp
View File

@ -637,10 +637,7 @@ void CompilerHLSL::emit_specialization_constants()
if (!c.specialization)
continue;
if (c.self == workgroup_size_id ||
c.self == wg_x.id ||
c.self == wg_y.id ||
c.self == wg_z.id)
if (c.self == workgroup_size_id || c.self == wg_x.id || c.self == wg_y.id || c.self == wg_z.id)
{
continue;
}
@ -689,9 +686,12 @@ void CompilerHLSL::emit_resources()
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (var.storage != StorageClassFunction && type.pointer && type.storage == StorageClassUniform &&
!is_hidden_variable(var) && (meta[type.self].decoration.decoration_flags &
((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))))
bool is_block_storage = type.storage == StorageClassStorageBuffer || type.storage == StorageClassUniform;
bool has_block_flags = (meta[type.self].decoration.decoration_flags &
((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))) != 0;
if (var.storage != StorageClassFunction && type.pointer && is_block_storage && !is_hidden_variable(var) &&
has_block_flags)
{
emit_buffer_block(var);
emitted = true;
@ -1021,7 +1021,7 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
bool is_uav = has_decoration(type.self, DecorationBufferBlock);
bool is_uav = var.storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock);
if (is_uav)
{
@ -1384,9 +1384,9 @@ void CompilerHLSL::emit_fixup()
if (options.shader_model <= 30)
{
statement("gl_Position.x = gl_Position.x - gl_HalfPixel.x * "
"gl_Position.w;");
"gl_Position.w;");
statement("gl_Position.y = gl_Position.y + gl_HalfPixel.y * "
"gl_Position.w;");
"gl_Position.w;");
}
if (CompilerGLSL::options.vertex.flip_vert_y)
@ -1813,6 +1813,8 @@ string CompilerHLSL::to_resource_binding(const SPIRVariable &var)
else
space = "c"; // Constant buffers
}
else if (storage == StorageClassStorageBuffer)
space = "u"; // UAV
break;
}
@ -2092,7 +2094,7 @@ void CompilerHLSL::emit_access_chain(const Instruction &instruction)
bool need_byte_access_chain = false;
auto &type = expression_type(ops[2]);
const SPIRAccessChain *chain = nullptr;
if (has_decoration(type.self, DecorationBufferBlock))
if (type.storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock))
{
// If we are starting to poke into an SSBO, we are dealing with ByteAddressBuffers, and we need
// to emit SPIRAccessChain rather than a plain SPIRExpression.

22
spirv_msl.cpp
View File

@ -217,7 +217,8 @@ void CompilerMSL::extract_global_variables_from_functions()
{
auto &var = id.get<SPIRVariable>();
if (var.storage == StorageClassInput || var.storage == StorageClassUniform ||
var.storage == StorageClassUniformConstant || var.storage == StorageClassPushConstant)
var.storage == StorageClassUniformConstant || var.storage == StorageClassPushConstant ||
var.storage == StorageClassStorageBuffer)
{
global_var_ids.insert(var.self);
}
@ -1004,7 +1005,7 @@ void CompilerMSL::emit_resources()
if (var.storage != StorageClassFunction && type.pointer &&
(type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant ||
type.storage == StorageClassPushConstant) &&
type.storage == StorageClassPushConstant || type.storage == StorageClassStorageBuffer) &&
(has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock)) &&
!is_hidden_variable(var))
{
@ -2272,12 +2273,17 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
if ((type.basetype == SPIRType::Struct) &&
(type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant ||
type.storage == StorageClassPushConstant))
type.storage == StorageClassPushConstant || type.storage == StorageClassStorageBuffer))
{
return ((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0 &&
(meta[argument.self].decoration.decoration_flags & (1ull << DecorationNonWritable)) == 0) ?
"device" :
"constant";
if (type.storage == StorageClassStorageBuffer)
return "device";
else
{
return ((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0 &&
(meta[argument.self].decoration.decoration_flags & (1ull << DecorationNonWritable)) == 0) ?
"device" :
"constant";
}
}
return "thread";
@ -2324,7 +2330,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
uint32_t var_id = var.self;
if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
var.storage == StorageClassPushConstant) &&
var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) &&
!is_hidden_variable(var))
{
switch (type.basetype)