Cross-kernel influence in binding allocation

[kpet]

Clvk currently assumes that the highest binding number in a descriptor set is less than the number of resources that require a binding. This no longer seems to be true with recent versions of clspv (kpet/clvk#233). While accomodating this in clvk is trivial, I wanted to double check that this bit of clspv behaviour was expected before doing so. Specifically, it seems that in some cases bindings for one kernel are allocated after those for another kernel. Consider the following code:

kernel void k0(int v, local int* l, global int* b){}
kernel void k1(global int* b){}

which leads to the following descriptor map:

kernel_decl,k0
kernel,k0,arg,l,argOrdinal,1,argKind,local,arrayElemSize,4,arrayNumElemSpecId,3
kernel,k0,arg,b,argOrdinal,2,descriptorSet,0,binding,0,offset,0,argKind,buffer
kernel,k0,arg,v,argOrdinal,0,offset,0,argKind,pod_pushconstant,argSize,4
kernel_decl,k1
kernel,k1,arg,b,argOrdinal,0,descriptorSet,0,binding,2,offset,0,argKind,buffer
spec_constant,workgroup_size_x,spec_id,0
spec_constant,workgroup_size_y,spec_id,1
spec_constant,workgroup_size_z,spec_id,2

Generally, this seems to happen when there are POD and local arguments. I haven't looked into the details.

Is this expected behaviour? If yes, what is the rationale for this behaviour?

[alan-baker]

If the sample was rewritten as:

kernel void k0(global int* b, int v, local int* l){}
kernel void k1(global int* b){}

Clspv would reuse binding 0 for by k0 and k1 b arguments. Currently clspv only shares bindings when the argument index and types match. This allows code generation to generate a single variable for the descriptor set and binding pair. That's why you end up with binding 2.

There are improvements in this regard. For example, DirectResourceAccess attempts to leverage reused bindings to avoid variable pointers. It should be the other way around and AllocateDescriptors should consider which resources can be shared before assigning bindings (a sort of graph colouring problem). Any change here though would necessitate changes in the SPIR-V producer to generate variables differently. This isn't necessarily difficult, but it is not currently accounted for.

[kpet]

I see, makes sense. Thank you for the explanation. I'm amazed I'm only bumping into this now! I'll push a fix to clvk.

[rjodinchr]

@kpet, Can we close this issue now?

[kpet]

Maybe @alan-baker wants to keep it to track his suggested improvements. Maybe a new issue would be better. I'm fine with closing it if Alan's fine.

[alan-baker]

I've filed #1061.