Module: Mesa
Branch: master
Commit: 5b92392c483c8f15599ed48abc824e606e859b02
URL:
http://cgit.freedesktop.org/mesa/mesa/commit/?id=5b92392c483c8f15599ed48abc824e606e859b02
Author: Rhys Perry <pendingchao...@gmail.com>
Date: Mon Jul 27 14:48:12 2020 +0100
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 <pendingchao...@gmail.com>
Reviewed-by: Boris Brezillon <boris.brezil...@collabora.com>
Reviewed-by: Jason Ekstrand <ja...@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6090>
---
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;
_______________________________________________
mesa-commit mailing list
mesa-commit@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-commit
