Fix bounds checks on pointers

[Liamolucko]

Draft fix for gfx-rs/wgpu#4541; this is still a draft because I haven't fixed the issue with SPIR-V yet (and have had to disable some tests as a result), but I'm opening it as a draft to check if my fix for Metal is sensible.

Like I proposed in the issue, this PR generates extra local variables which we give pointers to whenever code attempts to obtain a pointer that's out of bounds. I named these 'out-of-bounds locals' (OOB locals for short). One of them is generated per type that an out-of-bounds pointer might end up being created for.

They're only initialised to 0 once at the beginning of the function where they're declared, and from there on any writes to them remain visible; this means they don't technically fully uphold the ReadZeroSkipWrite bounds check policy. Writes are still effectively skipped, but reads can then observe prior writes. While it would be possible to restore them to 0 after every use, it doesn't really seem worth it to me when they can't get reset during the duration of any functions they're passed to anyway.

The main bit of my implementation that I'm not quite sure about is that I added two extra variants to NameKey which allow getting the names of the OOB locals of a given type within a function or entry point. This seems reasonable to me, since they're just another thing that we need to look up the name of, but it is slightly out of place as the only thing that isn't just assigning names to existing parts of a Module. Let me know if you'd rather I implemented that another way.

I also put some of the non-Metal specific logic into proc::index, since it seems potentially useful for any other backends that want to add bounds checking. (I am curious why Metal seems to be the only backend that uses bounds checking right now; do all the other languages already do it automatically? Or has it just not been implemented yet?)

Also, I'm testing this by just adding some more stuff to the existing access.wgsl snapshot test; however, that meant I had to enable BoundsCheckPolicy::ReadZeroSkipWrite for it. Does it matter that BoundsCheckPolicy::Unchecked is now no longer being tested, and is there a good way to test both?

[Liamolucko]

Oops, I didn't try to compile the Metal output for getting pointers to vector elements and it doesn't seem like Metal (or HLSL) actually lets you do that. I'm not really sure how to implement that then; WGSL's supposed to let you do it.

There's also some other HLSL errors popping up in code I didn't touch, not sure what's going on there.

[fornwall]

There's also some other HLSL errors popping up in code I didn't touch, not sure what's going on there.

That is due to a recent DirectX shader compiler release, should be fixed by #2447.

[Liamolucko]

I've just disabled the tests for getting pointers to vector/matrix elements for now, since they're already broken regardless of whether bounds checks are enabled.

I do have an idea for how that could be worked around, in a similar fashion to this PR: the element you want to get a pointer to is first copied into a temporary variable, a pointer to that is given to the function, and the element is copied back afterwards:

metal::float4 vec = metal::float4(1.0, 2.0, 3.0, 4.0);
float scratch = vec.y;
takes_ptr(scratch);
vec.y = scratch;

Out-of-bounds locals should probably then be merged into the same thing as those temporary / scratch variables (e.g. they'd both use the same NameKey::{Function,EntryPoint}Scratch). However that's outside the scope of this PR.

[Liamolucko]

I was able to fix the panic in SPIR-V, but it turns out that SPIR-V doesn't actually let you pass pointers to array elements to functions in the first place, let alone vector elements. You can't even pass pointers to struct fields: only pointers to variables and function arguments are allowed (with the exception of arrays of samplers and images).

This restriction is partially lifted by the VariablePointers capability, but only for pointers with the StorageBuffer or Workgroup storage class.

So I haven't bothered committing that code, since the only case where the panic occurs is when you pass a pointer that was created with an access chain to a function, which as I just mentioned wouldn't be valid SPIR-V anyway.

The only solution I can see to this is the same as the solution for vector indexing in Metal/HLSL: copy the array element into a local (or global, if the address space isn't function), pass a pointer to that local to the function, and copy it back afterwards. @teoxoy, would that be an acceptable solution?

Oh yeah, I also realised while working on this that my fix for Metal only currently works for pointers with the function address space since it always declares locals; I'll fix that momentarily.

[Liamolucko]

Oh yeah, I also realised while working on this that my fix for Metal only currently works for pointers with the function address space since it always declares locals; I'll fix that momentarily.

Never mind, I didn't realise that the function and private address spaces are indistinguishable in Metal (they both map to the thread address space), and those are the only two address spaces WGSL allows pointers to be in anyway.

For some reason Naga seems to also allow pointers in the workgroup address space though. Is that intentional, and do I need to support it here?

[cwfitzgerald]

Hello, thank you for your PR against Naga!

As part of gfx-rs/wgpu#4231, we have moved development of Naga into the wgpu repository in the Naga subfolder. We have transferred all issues, but we are unable to automatically transfer PRs.

As such, please recreate your PR against the wgpu repository. We apologize for the inconvenience this causes, but will make contributing to both projects more streamlined going forward.

We are leaving PRs open, but once they are transferred, please close the original Naga PR.