[Mesa-dev] [PATCH v3 057/104] nir,spirv: Rework function calls
Rob Clark
robdclark at gmail.com
Tue Apr 10 16:59:29 UTC 2018
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.)
>>
>> 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.
> 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.
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;
}
BR,
-R
More information about the mesa-dev
mailing list