This is a major reworking of this transform. The old transform code
was getting unwieldy, with part of the complication coming from the
handling of multiple return statements. By generating a wrapper
function instead, we can avoid a lot of this complexity.
The original entry point function is stripped of all shader IO
attributes (as well as `stage` and `workgroup_size`), but the body is
left unmodified. A new entry point wrapper function is introduced
which calls the original function, packing/unpacking the shader inputs
as necessary, and propagates the result to the corresponding shader
outputs.
The new code has been refactored to use a state object with the
different parts of the transform split into separate functions, which
makes it much more manageable.
Fixed: tint:1076
Bug: tint:920
Change-Id: I3490a0ea7a3509a4e198ce730e476516649d8d96
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60521
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Arrays can be extremely large, and having the load and store functions unroll the elements can make the complier explode.
Fixed: chromium:1229233
Change-Id: Ieb5654254e16f5ce724a205d21d954ef9a0cd053
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58382
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
CloneContext::Replace(T* what, T* with) is bug-prone, as complex transforms may want to clone `what` multiple times, or not at all. In both cases, this will likely result in an ICE as either the replacement will be reachable multiple times, or not at all.
The CTS test: webgpu:shader,execution,robust_access:linear_memory:storageClass="storage";storageMode="read_write";access="read";atomic=true;baseType="i32"
Was triggering this brokenness with DecomposeMemoryAccess's use of CloneContext::Replace(T*, T*).
Switch the usage of CloneContext::Replace(T*, T*) to the new function form.
As std::function is copyable, it cannot hold a captured std::unique_ptr.
This prevented the Replace() lambdas from capturing the necessary `BufferAccess` data, as this held a `std::unique_ptr<Offset>`.
To fix this, use a `BlockAllocator` for Offsets, and use raw pointers instead.
Because the function passed to Replace() is called just before the node is cloned, insertion of new functions will occur just before the currently evaluated module-scope entity.
This allows us to remove the "insert_after" arguments to LoadFunc(), StoreFunc(), and AtomicFunc().
We can also kill the icky InsertGlobal() and TypeDeclOf() helpers.
Bug: tint:993
Change-Id: I60972bc13a2fa819a163ee2671f61e82d0e68d2a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58222
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Use the new semantic constant value information to significantly reduce the complex indexing logic emitted for UBO accesses.
This will dramatically reduce the number of `for` loops that are decayed to `while` loops.
Change-Id: I1b0adb5edde2b4ed39c6beafc2e28106b86e0edd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57701
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Use the new transforms to try and simplify loops into for-loops.
Emit loops when the initialiser, condition and continuing are simple enough to do so.
Bug: tint:952
Change-Id: I5b3c225b245ffa72996abf6a70f52a9cd25b748e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56772
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
These indices were a mix of signed and unsigned.
Modulus on the signed integers was producing FXC warnings about performance.
Change-Id: Ib82f4296199a09d2f03be8b06314feefce0022e2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56765
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
For structures and arrays.
This behaves identically to the per-element zero-initialization, but can be significantly less verbose.
Change-Id: I380ef86f16c2b3f37a9de2820e707f368955b761
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56764
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
James Price
Ben Clayton