diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 2960acc7f71..2586b9ecba6 100644
--- 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.
+    * shaders of the same stage.  Global variables are declared per-module.
     *
-    * - 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.
-    *
-    * 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 7a7295b0363..28f6a80e1ab 100644
--- 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 39d0711b867..c21ac8194d8 100644
--- 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);