Commit ff02c8ca authored by Rhys Perry's avatar Rhys Perry Committed by Erik Faye-Lund
Browse files

spirv: fix Uniform and Output MemoryAccessMakePointer{Visible,Available}



The Uniform storage class can be used for SSBOs. This should also fix make
available/visible for the Output storage class.
Signed-off-by: Rhys Perry's avatarRhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Boris Brezillon's avatarBoris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand's avatarJason Ekstrand <jason@jlekstrand.net>
Part-of: <mesa/mesa!6090>
parent dd07bf0b
......@@ -2028,22 +2028,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
}
SpvMemorySemanticsMask
vtn_storage_class_to_memory_semantics(SpvStorageClass sc)
{
switch (sc) {
case SpvStorageClassStorageBuffer:
case SpvStorageClassPhysicalStorageBuffer:
return SpvMemorySemanticsUniformMemoryMask;
case SpvStorageClassWorkgroup:
return SpvMemorySemanticsWorkgroupMemoryMask;
case SpvStorageClassCrossWorkgroup:
return SpvMemorySemanticsCrossWorkgroupMemoryMask;
default:
return SpvMemorySemanticsMaskNone;
}
}
static void
vtn_split_barrier_semantics(struct vtn_builder *b,
SpvMemorySemanticsMask semantics,
......@@ -3513,7 +3497,7 @@ vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode,
/* Atomic ordering operations will implicitly apply to the atomic operation
* storage class, so include that too.
*/
semantics |= vtn_storage_class_to_memory_semantics(ptr->ptr_type->storage_class);
semantics |= vtn_mode_to_memory_semantics(ptr->mode);
SpvMemorySemanticsMask before_semantics;
SpvMemorySemanticsMask after_semantics;
......
......@@ -938,7 +938,7 @@ bool vtn_handle_amd_shader_explicit_vertex_parameter_instruction(struct vtn_buil
const uint32_t *words,
unsigned count);
SpvMemorySemanticsMask vtn_storage_class_to_memory_semantics(SpvStorageClass sc);
SpvMemorySemanticsMask vtn_mode_to_memory_semantics(enum vtn_variable_mode mode);
void vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
SpvMemorySemanticsMask semantics);
......
......@@ -2461,6 +2461,28 @@ vtn_get_mem_operands(struct vtn_builder *b, const uint32_t *w, unsigned count,
return true;
}
SpvMemorySemanticsMask
vtn_mode_to_memory_semantics(enum vtn_variable_mode mode)
{
switch (mode) {
case vtn_variable_mode_ssbo:
case vtn_variable_mode_phys_ssbo:
return SpvMemorySemanticsUniformMemoryMask;
case vtn_variable_mode_workgroup:
return SpvMemorySemanticsWorkgroupMemoryMask;
case vtn_variable_mode_cross_workgroup:
return SpvMemorySemanticsCrossWorkgroupMemoryMask;
case vtn_variable_mode_atomic_counter:
return SpvMemorySemanticsAtomicCounterMemoryMask;
case vtn_variable_mode_image:
return SpvMemorySemanticsImageMemoryMask;
case vtn_variable_mode_output:
return SpvMemorySemanticsOutputMemoryMask;
default:
return SpvMemorySemanticsMaskNone;
}
}
void
vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count)
......@@ -2579,7 +2601,7 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
if (access & SpvMemoryAccessMakePointerVisibleMask) {
SpvMemorySemanticsMask semantics =
SpvMemorySemanticsMakeVisibleMask |
vtn_storage_class_to_memory_semantics(src->ptr_type->storage_class);
vtn_mode_to_memory_semantics(src->mode);
vtn_emit_memory_barrier(b, scope, semantics);
}
......@@ -2628,7 +2650,7 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
if (access & SpvMemoryAccessMakePointerAvailableMask) {
SpvMemorySemanticsMask semantics =
SpvMemorySemanticsMakeAvailableMask |
vtn_storage_class_to_memory_semantics(dest->ptr_type->storage_class);
vtn_mode_to_memory_semantics(dest->mode);
vtn_emit_memory_barrier(b, scope, semantics);
}
break;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment