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
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);
}
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;
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_ */
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);