From: Alejandro Piñeiro <apinhe...@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.
---
src/compiler/spirv/nir_spirv.h | 5 ++
src/compiler/spirv/spirv_to_nir.c | 156 +++++++++++++++++++++++++++++++++++---
2 files changed, 151 insertions(+), 10 deletions(-)
diff --git a/src/compiler/spirv/nir_spirv.h b/src/compiler/spirv/nir_spirv.h
index d975254047a..f3ca26e590c 100644
--- a/src/compiler/spirv/nir_spirv.h
+++ b/src/compiler/spirv/nir_spirv.h
@@ -40,6 +40,7 @@ struct nir_spirv_specialization {
uint32_t data32;
uint64_t data64;
};
+ bool defined_on_module;
};
struct nir_spirv_supported_capabilities {
@@ -54,6 +55,10 @@ struct nir_spirv_supported_capabilities {
bool variable_pointers;
};
+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 6034228ed36..fea3ad7b09d 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1031,6 +1031,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;
}
}
@@ -1065,6 +1066,11 @@ 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;
+
assert(member == -1);
if (dec->decoration != SpvDecorationBuiltIn ||
dec->literals[0] != SpvBuiltInWorkgroupSize)
@@ -2848,6 +2854,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)
@@ -3056,6 +3105,22 @@ vtn_handle_variable_or_type_instruction(struct
vtn_builder *b, SpvOp opcode,
}
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)
{
@@ -3309,15 +3374,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 nir_spirv_supported_capabilities *cap,
- const nir_shader_compiler_options *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 uint32_t *word_end = words + word_count;
-
/* Handle the SPIR-V header (first 4 dwords) */
assert(word_count > 5);
@@ -3327,8 +3387,6 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
unsigned value_id_bound = words[3];
assert(words[4] == 0);
- words+= 5;
-
/* Initialize the stn_builder object */
struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
b->value_id_bound = value_id_bound;
@@ -3336,6 +3394,84 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
exec_list_make_empty(&b->functions);
b->entry_point_stage = stage;
b->entry_point_name = entry_point_name;
+
+ 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)
+{
+ const uint32_t *word_end = words + word_count;
+
+ struct vtn_builder *b = common_initialization(words, word_count,
+ stage, entry_point_name);
+ 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 nir_spirv_supported_capabilities *cap,
+ const nir_shader_compiler_options *options)
+{
+ const uint32_t *word_end = words + word_count;
+
+ struct vtn_builder *b = common_initialization(words, word_count,
+ stage, entry_point_name);
+ words+= 5;
b->cap = cap;
/* Handle all the preamble instructions */
--
2.11.0
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
