Attaching pipeline layout to hal.interface.binding.subspan & co.

[benvanik]

This allows for the whole layout to be known locally when lowering out of HAL and into target-specific binding data structures. This information was (and is still) available on the exports but annoying to get to and not present in all tests. This allowed removing the descriptor type from the subspan op and will allow for us to have non-i32 push constant types in the future. Verifiers were added for both push constant and descriptor set/binding ordinals now that the information is cheap to verify.

Progress on #17875 (this is needed for lowering non-0 ordinal descriptor sets to CPU/CUDA/ROCM targets).

[benvanik]

I split the codegen changes out from the actual dialect changes. The current failures are WGSL-related due to push constant emulation adding a descriptor set without updating the pipeline layout - I'll see if I can fix that. WGSL fixed!

[ScottTodd]

Nice having these explicit. Needing to look them up by walking the IR to find an export op was annoying :P

Can we ever end up using multiple pipeline layouts at this point? Maybe we could have a verifier that enforces only one?

#pipeline_layout_0 = #hal.pipeline.layout<push_constants = 4, sets = [
  #hal.descriptor_set.layout<0, bindings = [
    #hal.descriptor_set.binding<0, storage_buffer>
  ]>
]>

#pipeline_layout_1 = #hal.pipeline.layout<push_constants = 2, sets = [
  #hal.descriptor_set.layout<0, bindings = [
    #hal.descriptor_set.binding<0, storage_buffer>
  ]>
]>

hal.executable.source public @executable {
  hal.executable.export public @write_push_constants ordinal(0) layout(#pipeline_layout) attributes {workgroup_size = [1 : index, 1 : index, 1 : index]} {
  ^bb0(%arg0: !hal.device):
    %c1 = arith.constant 1 : index
    hal.return %c1, %c1, %c1 : index, index, index
  }
  builtin.module {
    func.func @write_push_constants() {
      %input_0 = hal.interface.constant.load layout(#pipeline_layout_0) ordinal(0) : i32
      %input_1 = hal.interface.constant.load layout(#pipeline_layout_1) ordinal(0) : i32

[benvanik]

It's possible so long as they are compatible - this could let us have multiple exported entry points that share some subset of buffers but have differing push constant counts, etc (so two top level functions that fetch push constants/buffers, then a common function that can still reference the base push constants/buffers). We could have a verifier on the executable that maybe poked in and looked for them, but it's likely going to need to be a dedicated pass given how much of the IR tree it pulls together. We'd also then have to codify what pipeline compatibility is and I'm not yet sure we know if we even want to allow that so I'm punting on it for now :P

[ScottTodd]

Yeah I figured the analysis would be complex enough that it would need to be a full pass and not a local verifier.