On 15/01/18 17:28, Jason Ekstrand wrote:
> On January 15, 2018 06:46:13 Alejandro Piñeiro <apinhe...@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.
Ok, makes sense. I just reused them in order to add as less code as
possible.
>
>> +
>> 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?
Ok.
>
>> {
>> /* 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.
Ok, will try that. If I find any other reason to keep the new method on
this file I will come back.
>
>> +{
>> + /* 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.
Ups. Sorry.
>
>> +
>> + 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.
Ups again.
>
>> +
>> + 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@lists.freedesktop.org
>> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
>
>
>
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
