From 5b92392c483c8f15599ed48abc824e606e859b02 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Mon, 27 Jul 2020 14:48:12 +0100 Subject: [PATCH] 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 Reviewed-by: Boris Brezillon Reviewed-by: Jason Ekstrand Part-of: --- src/compiler/spirv/spirv_to_nir.c | 16 +--------------- src/compiler/spirv/vtn_private.h | 2 +- src/compiler/spirv/vtn_variables.c | 26 ++++++++++++++++++++++++-- 3 files changed, 26 insertions(+), 18 deletions(-) diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 7324e9ca70b..8aa9ac0e34f 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -2103,20 +2103,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; - default: - return SpvMemorySemanticsMaskNone; - } -} - static void vtn_split_barrier_semantics(struct vtn_builder *b, SpvMemorySemanticsMask semantics, @@ -3539,7 +3525,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; diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h index 6fc4f2e4f3e..9c96f1dd2be 100644 --- a/src/compiler/spirv/vtn_private.h +++ b/src/compiler/spirv/vtn_private.h @@ -927,7 +927,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); diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c index 712c3db7415..fa78f804594 100644 --- a/src/compiler/spirv/vtn_variables.c +++ b/src/compiler/spirv/vtn_variables.c @@ -2478,6 +2478,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) @@ -2598,7 +2620,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); } @@ -2647,7 +2669,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; -- 2.30.2