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 <pendingchaos02@gmail.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6090>

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;

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);

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;