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