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

Merge pull request #266 from KhronosGroup/fix-265

Browse Source 
Support coherent SSBO
This commit is contained in:
Hans-Kristian Arntzen 2017-08-28 09:15:40 +02:00 committed by GitHub
commit 579a056e33
6 changed files with 48 additions and 25 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
reference/shaders/comp/coherent-block.comp Normal file
View File

@ -0,0 +1,13 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 1, std430) coherent restrict writeonly buffer SSBO
{
vec4 value;
} _10;
void main()
{
_10.value = vec4(20.0);
}

12
shaders/comp/coherent-block.comp Normal file
View File

@ -0,0 +1,12 @@
#version 310 es
layout(local_size_x = 1) in;
layout(binding = 1) coherent restrict writeonly buffer SSBO
{
vec4 value;
};
void main()
{
value = vec4(20.0);
}

14
spirv_common.hpp
View File

@ -614,13 +614,13 @@ struct SPIRAccessChain : IVariant
type = TypeAccessChain
};
SPIRAccessChain(uint32_t basetype_, spv::StorageClass storage_,
std::string base_, std::string dynamic_index_, int32_t static_index_)
: basetype(basetype_),
storage(storage_),
base(base_),
dynamic_index(std::move(dynamic_index_)),
static_index(static_index_)
SPIRAccessChain(uint32_t basetype_, spv::StorageClass storage_, std::string base_, std::string dynamic_index_,
int32_t static_index_)
: basetype(basetype_)
, storage(storage_)
, base(base_)
, dynamic_index(std::move(dynamic_index_))
, static_index(static_index_)
{
}

12
spirv_glsl.cpp
View File

@ -1113,6 +1113,7 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
bool is_restrict = ssbo && (flags & (1ull << DecorationRestrict)) != 0;
bool is_writeonly = ssbo && (flags & (1ull << DecorationNonReadable)) != 0;
bool is_readonly = ssbo && (flags & (1ull << DecorationNonWritable)) != 0;
bool is_coherent = ssbo && (flags & (1ull << DecorationCoherent)) != 0;
add_resource_name(var.self);
@ -1126,8 +1127,9 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
else
resource_names.insert(buffer_name);
statement(layout_for_variable(var), is_restrict ? "restrict " : "", is_writeonly ? "writeonly " : "",
is_readonly ? "readonly " : "", ssbo ? "buffer " : "uniform ", buffer_name);
statement(layout_for_variable(var), is_coherent ? "coherent " : "", is_restrict ? "restrict " : "",
is_writeonly ? "writeonly " : "", is_readonly ? "readonly " : "", ssbo ? "buffer " : "uniform ",
buffer_name);
begin_scope();
@ -4060,9 +4062,9 @@ std::string CompilerGLSL::flattened_access_chain_vector(uint32_t base, const uin
}
}
std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(const SPIRType &basetype, const uint32_t *indices,
uint32_t count, uint32_t offset,
uint32_t word_stride,
std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(const SPIRType &basetype,
const uint32_t *indices, uint32_t count,
uint32_t offset, uint32_t word_stride,
bool *need_transpose,
uint32_t *out_matrix_stride)
{

3
spirv_glsl.hpp
View File

@ -361,8 +361,7 @@ protected:
bool need_transpose);
std::pair<std::string, uint32_t> flattened_access_chain_offset(const SPIRType &basetype, const uint32_t *indices,
uint32_t count, uint32_t offset,
uint32_t word_stride,
bool *need_transpose = nullptr,
uint32_t word_stride, bool *need_transpose = nullptr,
uint32_t *matrix_stride = nullptr);
const char *index_to_swizzle(uint32_t index);

19
spirv_hlsl.cpp
View File

@ -956,8 +956,8 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
uint64_t flags = get_buffer_block_flags(var);
bool is_readonly = (flags & (1ull << DecorationNonWritable)) != 0;
add_resource_name(var.self);
statement(is_readonly ? "ByteAddressBuffer " : "RWByteAddressBuffer ",
to_name(var.self), type_to_array_glsl(type), to_resource_binding(var), ";");
statement(is_readonly ? "ByteAddressBuffer " : "RWByteAddressBuffer ", to_name(var.self),
type_to_array_glsl(type), to_resource_binding(var), ";");
}
else
{
@ -1084,8 +1084,8 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, uint64_t return_f
{
// Manufacture automatic sampler arg for SampledImage texture
decl += ", ";
decl += join(arg_type.image.depth ? "SamplerComparisonState " : "SamplerState ",
to_sampler_expression(arg.id));
decl +=
join(arg_type.image.depth ? "SamplerComparisonState " : "SamplerState ", to_sampler_expression(arg.id));
}
if (&arg != &func.arguments.back())
@ -1138,9 +1138,7 @@ void CompilerHLSL::emit_hlsl_entry_point()
if (execution.model == ExecutionModelGLCompute)
{
statement("[numthreads(",
execution.workgroup_size.x, ", ",
execution.workgroup_size.y, ", ",
statement("[numthreads(", execution.workgroup_size.x, ", ", execution.workgroup_size.y, ", ",
execution.workgroup_size.z, ")]");
}
@ -2031,10 +2029,9 @@ void CompilerHLSL::emit_access_chain(const Instruction &instruction)
uint32_t matrix_stride = 0;
bool need_transpose = false;
auto offsets = flattened_access_chain_offset(*basetype,
&ops[3 + to_plain_buffer_length], length - 3 - to_plain_buffer_length,
0, 1, &need_transpose, &matrix_stride);
auto offsets =
flattened_access_chain_offset(*basetype, &ops[3 + to_plain_buffer_length],
length - 3 - to_plain_buffer_length, 0, 1, &need_transpose, &matrix_stride);
auto &e = set<SPIRAccessChain>(ops[1], ops[0], type.storage, base, offsets.first, offsets.second);
if (chain)