[Mesa-dev] [PATCH v3 03/10] nir/spirv: add gl_spirv_validation method

Eduardo Lima Mitev elima at igalia.com
Wed Dec 13 19:32:49 UTC 2017


From: Alejandro PiƱeiro <apinheiro at igalia.com>

ARB_gl_spirv adds the ability to use SPIR-V binaries, and a new
method, glSpecializeShader. From OpenGL 4.6 spec, section 7.2.1
"Shader Specialization", error table:

   INVALID_VALUE is generated if <pEntryPoint> does not name a valid
   entry point for <shader>.

   INVALID_VALUE is generated if any element of <pConstantIndex>
   refers to a specialization constant that does not exist in the
   shader module contained in <shader>.""

But we are not really interested on creating the nir shader at that
point, and adding nir structures on the gl_program, so at that point
we are just interested on the error checking.

So we add a new method focused on just checking those errors. It still
needs to parse the binary, but skips what it is not needed, and
doesn't create the nir shader.

v2: rebase update (spirv_to_nir options added, changes on the warning
    logging, and others)
v3: include passing options on common initialization, doesn't call
    setjmp on common_initialization
---
 src/compiler/spirv/nir_spirv.h    |   5 +
 src/compiler/spirv/spirv_to_nir.c | 191 ++++++++++++++++++++++++++++++++++----
 2 files changed, 180 insertions(+), 16 deletions(-)

diff --git a/src/compiler/spirv/nir_spirv.h b/src/compiler/spirv/nir_spirv.h
index a2c40e57d18..d2766abb7f9 100644
--- a/src/compiler/spirv/nir_spirv.h
+++ b/src/compiler/spirv/nir_spirv.h
@@ -41,6 +41,7 @@ struct nir_spirv_specialization {
       uint32_t data32;
       uint64_t data64;
    };
+   bool defined_on_module;
 };
 
 enum nir_spirv_debug_level {
@@ -69,6 +70,10 @@ struct spirv_to_nir_options {
    } debug;
 };
 
+bool gl_spirv_validation(const uint32_t *words, size_t word_count,
+                         struct nir_spirv_specialization *spec, unsigned num_spec,
+                         gl_shader_stage stage, const char *entry_point_name);
+
 nir_function *spirv_to_nir(const uint32_t *words, size_t word_count,
                            struct nir_spirv_specialization *specializations,
                            unsigned num_specializations,
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 0493dd3a8b1..7f959711be6 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1227,6 +1227,7 @@ spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
             const_value->data64 = b->specializations[i].data64;
          else
             const_value->data32 = b->specializations[i].data32;
+         b->specializations[i].defined_on_module = true;
          return;
       }
    }
@@ -1261,7 +1262,13 @@ handle_workgroup_size_decoration_cb(struct vtn_builder *b,
                                     const struct vtn_decoration *dec,
                                     void *data)
 {
+   /* This can happens if we are gl_spirv_validation. We can return safely, as
+    * we don't need the workgroup info for such validation. */
+   if (b->shader == NULL)
+      return;
+
    vtn_assert(member == -1);
+
    if (dec->decoration != SpvDecorationBuiltIn ||
        dec->literals[0] != SpvBuiltInWorkgroupSize)
       return;
@@ -3162,6 +3169,49 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
    return true;
 }
 
+/*
+ * gl_spirv validation. Just need to check for the entry point.
+ */
+static bool
+vtn_validate_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
+                                  const uint32_t *w, unsigned count)
+{
+   switch (opcode) {
+   /* The following opcodes are not needed for gl_spirv, so we can skip
+    * them.
+    */
+   case SpvOpSource:
+   case SpvOpSourceExtension:
+   case SpvOpSourceContinued:
+   case SpvOpExtension:
+   case SpvOpCapability:
+   case SpvOpExtInstImport:
+   case SpvOpMemoryModel:
+   case SpvOpString:
+   case SpvOpName:
+   case SpvOpMemberName:
+   case SpvOpExecutionMode:
+   case SpvOpDecorationGroup:
+   case SpvOpMemberDecorate:
+   case SpvOpGroupDecorate:
+   case SpvOpGroupMemberDecorate:
+      break;
+
+   case SpvOpEntryPoint:
+      vtn_handle_preamble_instruction(b, opcode, w, count);
+      break;
+
+   case SpvOpDecorate:
+      vtn_handle_decoration(b, opcode, w, count);
+      break;
+
+   default:
+      return false; /* End of preamble */
+   }
+
+   return true;
+}
+
 static void
 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
                           const struct vtn_decoration *mode, void *data)
@@ -3371,6 +3421,22 @@ vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
    return true;
 }
 
+static bool
+vtn_handle_constant_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
+                                        const uint32_t *w, unsigned count)
+{
+   switch (opcode) {
+   case SpvOpUndef:
+   case SpvOpVariable:
+      break;
+
+   default:
+      return vtn_handle_variable_or_type_instruction(b, opcode, w, count);
+   }
+
+   return true;
+}
+
 static bool
 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
                             const uint32_t *w, unsigned count)
@@ -3657,12 +3723,10 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
    return true;
 }
 
-nir_function *
-spirv_to_nir(const uint32_t *words, size_t word_count,
-             struct nir_spirv_specialization *spec, unsigned num_spec,
-             gl_shader_stage stage, const char *entry_point_name,
-             const struct spirv_to_nir_options *options,
-             const nir_shader_compiler_options *nir_options)
+static struct vtn_builder*
+common_initialization(const uint32_t *words, size_t word_count,
+                      gl_shader_stage stage, const char *entry_point_name,
+                      const struct spirv_to_nir_options *options)
 {
    /* Initialize the stn_builder object */
    struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
@@ -3675,14 +3739,6 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    b->entry_point_name = entry_point_name;
    b->options = options;
 
-   /* See also _vtn_fail() */
-   if (setjmp(b->fail_jump)) {
-      ralloc_free(b);
-      return NULL;
-   }
-
-   const uint32_t *word_end = words + word_count;
-
    /* Handle the SPIR-V header (first 4 dwords)  */
    vtn_assert(word_count > 5);
 
@@ -3692,11 +3748,114 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    unsigned value_id_bound = words[3];
    vtn_assert(words[4] == 0);
 
-   words+= 5;
-
    b->value_id_bound = value_id_bound;
    b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
 
+   return b;
+}
+
+/*
+ * Since OpenGL 4.6 you can use SPIR-V modules directly on OpenGL. One of the
+ * new methods, glSpecializeShader include some possible errors when trying to
+ * use it. From OpenGL 4.6, Section 7.2.1, "Shader Specialization":
+ *
+ * "void SpecializeShaderARB(uint shader,
+ *                           const char* pEntryPoint,
+ *                           uint numSpecializationConstants,
+ *                           const uint* pConstantIndex,
+ *                           const uint* pConstantVaulue);
+ * <skip>
+ *
+ * INVALID_VALUE is generated if <pEntryPoint> does not name a valid
+ * entry point for <shader>.
+ *
+ * An INVALID_VALUE error is generated if any element of pConstantIndex refers
+ * to a specialization constant that does not exist in the shader module
+ * contained in shader."
+ *
+ * We could do those checks on spirv_to_nir, but we are only interested on the
+ * full translation later, during linking. This method is a simplified version
+ * of spirv_to_nir, looking for only the checks needed by SpecializeShader.
+ *
+ * This method returns NULL if no entry point was found, and fill the
+ * nir_spirv_specialization field "defined_on_module" accordingly. Caller
+ * would need to trigger the specific errors.
+ *
+ */
+bool
+gl_spirv_validation(const uint32_t *words, size_t word_count,
+                    struct nir_spirv_specialization *spec, unsigned num_spec,
+                    gl_shader_stage stage, const char *entry_point_name)
+{
+   /* vtn_warn/vtn_log uses debug.func. Setting a null to prevent crash. Not
+    * need to print the warnings now, would be done later, on the real
+    * spirv_to_nir
+    */
+   const struct spirv_to_nir_options options = { .debug.func = NULL};
+   const uint32_t *word_end = words + word_count;
+
+   struct vtn_builder *b = common_initialization(words, word_count,
+                                                 stage, entry_point_name,
+                                                 &options);
+
+   /* See also _vtn_fail() */
+   if (setjmp(b->fail_jump)) {
+      ralloc_free(b);
+      return false;
+   }
+
+   if (b == NULL)
+      return false;
+
+   words+= 5;
+
+   /* Search entry point from preamble */
+   words = vtn_foreach_instruction(b, words, word_end,
+                                   vtn_validate_preamble_instruction);
+
+   if (b->entry_point == NULL) {
+      ralloc_free(b);
+      return false;
+   }
+
+   b->specializations = spec;
+   b->num_specializations = num_spec;
+
+   /* Handle type, and constant instructions (we don't need to handle
+    * variables for gl_spirv)
+    */
+   words = vtn_foreach_instruction(b, words, word_end,
+                                   vtn_handle_constant_or_type_instruction);
+
+   ralloc_free(b);
+
+   return true;
+}
+
+nir_function *
+spirv_to_nir(const uint32_t *words, size_t word_count,
+             struct nir_spirv_specialization *spec, unsigned num_spec,
+             gl_shader_stage stage, const char *entry_point_name,
+             const struct spirv_to_nir_options *options,
+             const nir_shader_compiler_options *nir_options)
+
+{
+   const uint32_t *word_end = words + word_count;
+
+   struct vtn_builder *b = common_initialization(words, word_count,
+                                                 stage, entry_point_name,
+                                                 options);
+   /* See also _vtn_fail() */
+   if (setjmp(b->fail_jump)) {
+      ralloc_free(b);
+      return NULL;
+   }
+
+   if (b == NULL)
+      return NULL;
+
+   words+= 5;
+
    /* Handle all the preamble instructions */
    words = vtn_foreach_instruction(b, words, word_end,
                                    vtn_handle_preamble_instruction);
-- 
2.15.1



More information about the mesa-dev mailing list