[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