[Mesa-dev] [PATCH v3 057/104] nir,spirv: Rework function calls
Jason Ekstrand
jason at jlekstrand.net
Tue Apr 10 17:17:46 UTC 2018
On Tue, Apr 10, 2018 at 9:59 AM, Rob Clark <robdclark at gmail.com> wrote:
> On Tue, Apr 10, 2018 at 11:55 AM, Jason Ekstrand <jason at jlekstrand.net>
> wrote:
> > On Tue, Apr 10, 2018 at 8:17 AM, Rob Clark <robdclark at gmail.com> wrote:
> >>
> >> On Tue, Apr 10, 2018 at 11:04 AM, Jason Ekstrand <jason at jlekstrand.net>
> >> wrote:
> >> > On Tue, Apr 10, 2018 at 6:20 AM, Rob Clark <robdclark at gmail.com>
> wrote:
> >> >>
> >> >> On Mon, Apr 9, 2018 at 10:52 PM, Jason Ekstrand <
> jason at jlekstrand.net>
> >> >> wrote:
> >> >> > + A bunch of potentially interested parties.
> >> >> >
> >> >> > On Mon, Apr 9, 2018 at 4:25 PM, Caio Marcelo de Oliveira Filho
> >> >> > <caio.oliveira at intel.com> wrote:
> >> >> >>
> >> >> >> Hi,
> >> >> >>
> >> >> >> > typedef struct {
> >> >> >> > - nir_parameter_type param_type;
> >> >> >> > - const struct glsl_type *type;
> >> >> >> > + uint8_t num_components;
> >> >> >> > + uint8_t bit_size;
> >> >> >> > } nir_parameter;
> >> >> >>
> >> >> >> (...)
> >> >> >>
> >> >> >> > @@ -683,18 +692,12 @@ validate_tex_instr(nir_tex_instr *instr,
> >> >> >> > validate_state *state)
> >> >> >> > static void
> >> >> >> > validate_call_instr(nir_call_instr *instr, validate_state
> *state)
> >> >> >> > {
> >> >> >> > - if (instr->return_deref == NULL) {
> >> >> >> > - validate_assert(state,
> >> >> >> > glsl_type_is_void(instr->callee->return_type));
> >> >> >> > - } else {
> >> >> >> > - validate_assert(state, instr->return_deref->deref.type
> ==
> >> >> >> > instr->callee->return_type);
> >> >> >> > - validate_deref_var(instr, instr->return_deref, state);
> >> >> >> > - }
> >> >> >> > -
> >> >> >> > validate_assert(state, instr->num_params ==
> >> >> >> > instr->callee->num_params);
> >> >> >> >
> >> >> >> > for (unsigned i = 0; i < instr->num_params; i++) {
> >> >> >> > - validate_assert(state, instr->callee->params[i].type ==
> >> >> >> > instr->params[i]->deref.type);
> >> >> >> > - validate_deref_var(instr, instr->params[i], state);
> >> >> >> > + validate_src(&instr->params[i], state,
> >> >> >> > + instr->callee->params[i].bit_size,
> >> >> >> > + instr->callee->params[i].num_components);
> >> >> >> > }
> >> >> >> > }
> >> >> >>
> >> >> >> Question: I might be misreading, but it seems like we are losing
> the
> >> >> >> type information for functions. Isn't that something worth
> keeping,
> >> >> >> maybe in some other way, e.g. load_param specifying the expected
> >> >> >> type?
> >> >> >
> >> >> >
> >> >> > That's a very good question! To be honest, I'm not sure what the
> >> >> > answer
> >> >> > is.
> >> >> > At the moment, the type information is fairly useless for most of
> >> >> > what
> >> >> > we
> >> >> > use functions for. Really, all we need is something that NIR can
> >> >> > inline.
> >> >> > As it is, we're not really preserving the types from SPIR-V because
> >> >> > of
> >> >> > the
> >> >> > gymnastics we're doing to handle pointers.
> >> >> >
> >> >> > If we did want to preserve types, we'd need to have more detailed
> >> >> > type
> >> >> > information. In particular, we'd need to be able to provide
> pointer
> >> >> > types
> >> >> > and maybe combined image-sampler types. And along with those
> pointer
> >> >> > types,
> >> >> > we'd need to somehow express those pointer's storage requirements.
> >> >> >
> >> >> > The philosophy behind this commit is that, if we don't have a good
> >> >> > match
> >> >> > to
> >> >> > SPIR-V anyway, we might as well just chuck that information and do
> >> >> > whatever
> >> >> > makes our lives the easiest. My philosophy here may be flawed and
> >> >> > I'm
> >> >> > happy
> >> >> > to hear arguments in favor of keeping the information. The best
> >> >> > argument I
> >> >> > can come up with for keeping the information is if we find
> ourselves
> >> >> > wanting
> >> >> > to do some sort of linking in the future where we have to match
> >> >> > functions by
> >> >> > both name and type. If we want to do that, however, we'll need all
> >> >> > the
> >> >> > SPIR-V type information.
> >> >> >
> >> >>
> >> >> We do end up wanting the type information for cl kernels. This is
> >> >> maybe a slightly different case from calls within shader code (ie.
> >> >> when both caller and callee are in shader).
> >> >
> >> >
> >> > Yes, I think it is. Question: Is there a distinction in CL between
> >> > functions which are entrypoints callable from the API and functions
> >> > which
> >> > are helpers? i.e. Can you call an entrypoint as a helper?
> >> >
> >>
> >> There is the __kernel annotation. And you know the entry point name
> >> when compiling. However I'm not sure anything prevents one entry
> >> point from calling another.
> >
> >
> > That would be worth investigating.
> >
>
> fwiw, at least the cl to spv compiler seems to allow it. (Although in
> my simple examples it also inlines the called function before things
> end up in spv.)
>
Interesting.
> >>
> >> I'm not sure we want the calling convention to be the same internally
> >> as for kernel entry points so in that case, if we aren't inlining
> >> everything, we might end up generating two versions of a function (or
> >> possibly a shim.. or possibly between the two based on size.. or??)
> >
> >
> > Having a shim seems like a reasonable plan.
> >
> >>
> >> >>
> >> >> Although I'd kinda like
> >> >> to think that we don't need to make vtn aware of this distinction.
> >> >
> >> >
> >> > Someone has to be aware of it. :-) There are lots of places in
> >> > spirv_to_nir
> >> > were we take the SPIR-V and do something slightly different with it
> than
> >> > the
> >> > obvious translation. Also, using function parameters for this is a
> >> > significant anachronism because no other shader I/O in NIR has ever
> >> > worked
> >> > that way.
> >> >
> >> >>
> >> >> So just to throw out an idea. What if vtn just used load_deref for
> >> >> everything, and in the case of fxn params it just points to a local
> >> >> var with type nir_var_param? (Or something roughly like that.) Then
> >> >> lower_io lowers this to load_param.
> >> >
> >> >
> >> > That's kind-of what the original thing did. However, for SPIR-V
> helper
> >> > functions we have to be able to pass through pointers, SSA values with
> >> > arbitrary type, and image/sampler pointers. SSA values can be handled
> >> > by
> >> > just making a variable and storing them to it. Pointers are tricky
> >> > because
> >> > they're not really copy-in/out. For images, samplers, and pointers,
> we
> >> > have
> >> > a pile of "try to patch up the deref chain" code in
> nir_inline_functions
> >> > that's rather tricky. The moral of the story is that "just use
> >> > variables"
> >> > is not nearly as obvious of a choice as it looks.
> >>
> >> So, I'm considering just adding nir_variable ptr back to
> >> nir_parameter. That would be rather easy, and gives something that
> >> some lowering pass could fixup data.driver_loc with parameter
> >> position. That would be rather easy for driver backends:
> >>
> >> param_loc = fxn->params[param_idx]->var->data.driver_loc
> >
> >
> > Why can't you just add a mode nir_var_kernel_param and keep a running
> > uint32_t offset in the builder. When you process an OpParameter in an
> > entrypoint, create a nir_var_kernel_param variable and set
> > var->data.location = offset and add the size of the param to offset.
> Then
>
> So far, with the new deref chain, I've managed to keep knowledge of
> physical layout of things in lower_io and out of vtn.. it would be
> nice to keep it that way. Otherwise vtn needs to know about
> size/alignment rules.
>
Another option (which maybe you've already brought up without me
understanding) would be to make nir_var_kernel_param a local variable type
and have the list of params stored in the function next to the locals.
Then you would know exactly what params are associated with a given
function impl and can assign locations at-will.
> > the OpParameter can return an access chain to the newly created interface
> > variable. I fail to see why load_param is needed at all here.
>
> I'd prefer not to have this end up like load_input (or load_uniform)
> in the driver (if that is what you had in mind).. *maybe* we could end
> up treating this like a UBO, although blob does try to use uniforms
> when possible (like when it doesn't need a pointer or byte
> addressing), since that is faster. So I think we should have a
> distinct intrinsic.. whether that be load_param or something
> different.
>
What is it about load_input and load_uniform that you don't like? I'm
really trying to understand what's going on here. Clearly, there's
something but I don't know what it is.
> Also, fwiw, we need to be able to do pointer math on function parameters
> ;-)
>
> __kernel void test_fn(__global int *res, __constant struct SomeStruct
> b[5], int c)
> {
> __constant struct SomeStruct *s = &b[c];
> if (c < 3)
> s++;
> *res = s->s.b;
> }
>
I figured you would. That's part of my confusion about why you want to
treat them the same as parameters you pass from one function to another.
They don't work at all like local variables.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180410/62fc78ea/attachment-0001.html>
More information about the mesa-dev
mailing list