spirv: Use OpEntryPoint to identify valid I/O variables

OpEntryPoint declares the list of variables in Input and Output
storage classes that are used.  Use that information to skip creating
other variables from such storage classes that are unused by the entry
point.

After that change, is not necessary to use remove dead variables for
those types of variables; and because of that is also not necessary to
lower initalizers for output variables.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8456>

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 2960acc..2586b9e 100644 (file)
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -4188,6 +4188,13 @@ vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
 
    vtn_assert(b->entry_point == NULL);
    b->entry_point = entry_point;
+
+   /* Entry points enumerate which I/O variables are used. */
+   size_t start = 3 + name_words;
+   b->interface_ids_count = count - start;
+   b->interface_ids = ralloc_array(b, uint32_t, b->interface_ids_count);
+   memcpy(b->interface_ids, &w[start], b->interface_ids_count * 4);
+   qsort(b->interface_ids, b->interface_ids_count, 4, cmp_uint32_t);
 }
 
 static bool
@@ -5966,27 +5973,21 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    nir_lower_goto_ifs(b->shader);
 
    /* A SPIR-V module can have multiple shaders stages and also multiple
-    * shaders of the same stage.  Global variables are declared per-module, so
-    * they are all collected when parsing a single shader.  These dead
-    * variables can result in invalid NIR, e.g.
-    *
-    * - TCS outputs must be per-vertex arrays (or decorated 'patch'), while VS
-    *   output variables wouldn't be;
-    * - Two vertex shaders have two different typed blocks associated to the
-    *   same Binding.
+    * shaders of the same stage.  Global variables are declared per-module.
     *
-    * Before cleaning the dead variables, we must lower any constant
-    * initializers on outputs so nir_remove_dead_variables sees that they're
-    * written to.
+    * For I/O storage classes, OpEntryPoint will list the variables used, so
+    * only valid ones are created.  Remove dead variables to clean up the
+    * remaining ones.
     */
    if (!options->create_library) {
-      nir_lower_variable_initializers(b->shader, nir_var_shader_out |
-                                                 nir_var_system_value);
       const nir_remove_dead_variables_options dead_opts = {
          .can_remove_var = can_remove,
          .can_remove_var_data = b->vars_used_indirectly,
       };
-      nir_remove_dead_variables(b->shader, ~nir_var_function_temp,
+      nir_remove_dead_variables(b->shader, ~(nir_var_function_temp |
+                                             nir_var_shader_out |
+                                             nir_var_shader_in |
+                                             nir_var_system_value),
                                 b->vars_used_indirectly ? &dead_opts : NULL);
    }
 

diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 7a7295b..28f6a80 100644 (file)
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -708,6 +708,9 @@ struct vtn_builder {
    struct vtn_value *workgroup_size_builtin;
    bool variable_pointers;
 
+   uint32_t *interface_ids;
+   size_t interface_ids_count;
+
    struct vtn_function *func;
    struct list_head functions;
 
@@ -1002,4 +1005,16 @@ SpvMemorySemanticsMask vtn_mode_to_memory_semantics(enum vtn_variable_mode mode)
 void vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
                              SpvMemorySemanticsMask semantics);
 
+static inline int
+cmp_uint32_t(const void *pa, const void *pb)
+{
+   uint32_t a = *((const uint32_t *)pa);
+   uint32_t b = *((const uint32_t *)pb);
+   if (a < b)
+      return -1;
+   if (a > b)
+      return 1;
+   return 0;
+}
+
 #endif /* _VTN_PRIVATE_H_ */

diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c
index 39d0711..c21ac81 100644 (file)
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -2291,9 +2291,16 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
    case SpvOpVariable: {
       struct vtn_type *ptr_type = vtn_get_type(b, w[1]);
 
-      struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_pointer);
-
       SpvStorageClass storage_class = w[3];
+
+      /* Skip I/O variables that are not used by the entry point. */
+      if (!b->options->create_library &&
+          (storage_class == SpvStorageClassInput ||
+           storage_class == SpvStorageClassOutput) &&
+          !bsearch(&w[2], b->interface_ids, b->interface_ids_count, 4, cmp_uint32_t))
+         break;
+
+      struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_pointer);
       struct vtn_value *initializer = count > 4 ? vtn_untyped_value(b, w[4]) : NULL;
 
       vtn_create_variable(b, val, ptr_type, storage_class, initializer);