[Mesa-dev] [PATCH v4 03/10] nir/spirv: add gl_spirv_validation method
Jason Ekstrand
jason at jlekstrand.net
Mon Jan 15 16:28:11 UTC 2018
On January 15, 2018 06:46:13 Alejandro Piñeiro <apinheiro at igalia.com> wrote:
> 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 c6df764682e..2143cd9df31 100644
> --- a/src/compiler/spirv/spirv_to_nir.c
> +++ b/src/compiler/spirv/spirv_to_nir.c
> @@ -1332,6 +1332,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;
> }
> }
> @@ -1366,7 +1367,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;
I don't think that re-using these two functions is really buying us
anything. We could just make spec constant validation versions that just
do what's needed there.
> +
> vtn_assert(member == -1);
> +
> if (dec->decoration != SpvDecorationBuiltIn ||
> dec->literals[0] != SpvBuiltInWorkgroupSize)
> return;
> @@ -3263,6 +3270,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)
> @@ -3473,6 +3523,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)
> {
> @@ -3775,12 +3841,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)
How about vtn_builder_create?
> {
> /* Initialize the stn_builder object */
> struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
> @@ -3794,14 +3858,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);
>
> @@ -3811,11 +3867,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)
Would it be reasonable to out this in it's own file? It seems to me like
the only thing you really need to re-use is handle_decoration and the other
attempts at code re-use are just confusing things.
> +{
> + /* 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;
These two checks are in the wrong order.
> +
> + 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;
Again, the null check needs to go first.
> +
> + words+= 5;
> +
> /* Handle all the preamble instructions */
> words = vtn_foreach_instruction(b, words, word_end,
> vtn_handle_preamble_instruction);
> --
> 2.11.0
>
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
More information about the mesa-dev
mailing list