Mesa (master): vtn: add an option to create a nir library from spirv

[GitLab Mirror]

Module: Mesa
Branch: master
Commit: c9a6b94f9ea50f73f016db4724472d87c9fb2334
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=c9a6b94f9ea50f73f016db4724472d87c9fb2334

Author: Dave Airlie <airlied at redhat.com>
Date:   Mon Nov 18 16:39:09 2019 +1000

vtn: add an option to create a nir library from spirv

This adds an options to turn a spir-v library into a set of nir
functions. There is no entry point and all function implentations
are emitted.

Reviewed-by: Jesse Natalie <jenatali at microsoft.com>
Reviewed-by: Jason Ekstrand <jason at jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6035>

---

 src/compiler/spirv/nir_spirv.h    |  3 +++
 src/compiler/spirv/spirv_to_nir.c | 38 ++++++++++++++++++++++----------------
 2 files changed, 25 insertions(+), 16 deletions(-)

diff --git a/src/compiler/spirv/nir_spirv.h b/src/compiler/spirv/nir_spirv.h
index 8d5c1008aa2..26f5024edc6 100644
--- a/src/compiler/spirv/nir_spirv.h
+++ b/src/compiler/spirv/nir_spirv.h
@@ -68,6 +68,9 @@ struct spirv_to_nir_options {
     */
    bool view_index_is_input;
 
+   /* Create a nir library. */
+   bool create_library;
+
    struct spirv_supported_capabilities caps;
 
    /* Address format for various kinds of pointers. */
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 4a6fcdfca6c..ea042bd57b6 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -5666,7 +5666,7 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    words = vtn_foreach_instruction(b, words, word_end,
                                    vtn_handle_preamble_instruction);
 
-   if (b->entry_point == NULL) {
+   if (!options->create_library && b->entry_point == NULL) {
       vtn_fail("Entry point not found for %s shader \"%s\"",
                _mesa_shader_stage_to_string(stage), entry_point_name);
       ralloc_free(b);
@@ -5682,8 +5682,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
       b->shader->info.gs.invocations = 1;
 
    /* Parse execution modes. */
-   vtn_foreach_execution_mode(b, b->entry_point,
-                              vtn_handle_execution_mode, NULL);
+   if (!options->create_library)
+      vtn_foreach_execution_mode(b, b->entry_point,
+                                 vtn_handle_execution_mode, NULL);
 
    b->specializations = spec;
    b->num_specializations = num_spec;
@@ -5695,8 +5696,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    /* Parse execution modes that depend on IDs. Must happen after we have
     * constants parsed.
     */
-   vtn_foreach_execution_mode(b, b->entry_point,
-                              vtn_handle_execution_mode_id, NULL);
+   if (!options->create_library)
+      vtn_foreach_execution_mode(b, b->entry_point,
+                                 vtn_handle_execution_mode_id, NULL);
 
    if (b->workgroup_size_builtin) {
       vtn_assert(b->workgroup_size_builtin->type->type ==
@@ -5715,15 +5717,17 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
 
    vtn_build_cfg(b, words, word_end);
 
-   assert(b->entry_point->value_type == vtn_value_type_function);
-   b->entry_point->func->referenced = true;
+   if (!options->create_library) {
+      assert(b->entry_point->value_type == vtn_value_type_function);
+      b->entry_point->func->referenced = true;
+   }
 
    bool progress;
    do {
       progress = false;
       vtn_foreach_cf_node(node, &b->functions) {
          struct vtn_function *func = vtn_cf_node_as_function(node);
-         if (func->referenced && !func->emitted) {
+         if ((options->create_library || func->referenced) && !func->emitted) {
             b->const_table = _mesa_pointer_hash_table_create(b);
 
             vtn_function_emit(b, func, vtn_handle_body_instruction);
@@ -5732,19 +5736,21 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
       }
    } while (progress);
 
-   vtn_assert(b->entry_point->value_type == vtn_value_type_function);
-   nir_function *entry_point = b->entry_point->func->impl->function;
-   vtn_assert(entry_point);
+   if (!options->create_library) {
+      vtn_assert(b->entry_point->value_type == vtn_value_type_function);
+      nir_function *entry_point = b->entry_point->func->impl->function;
+      vtn_assert(entry_point);
+
+      /* post process entry_points with input params */
+      if (entry_point->num_params && b->shader->info.stage == MESA_SHADER_KERNEL)
+         entry_point = vtn_emit_kernel_entry_point_wrapper(b, entry_point);
 
-   /* post process entry_points with input params */
-   if (entry_point->num_params && b->shader->info.stage == MESA_SHADER_KERNEL)
-      entry_point = vtn_emit_kernel_entry_point_wrapper(b, entry_point);
+      entry_point->is_entrypoint = true;
+   }
 
    /* structurize the CFG */
    nir_lower_goto_ifs(b->shader);
 
-   entry_point->is_entrypoint = true;
-
    /* When multiple shader stages exist in the same SPIR-V module, we
     * generate input and output variables for every stage, in the same
     * NIR program.  These dead variables can be invalid NIR.  For example,