Emit the "[loop]" attribute on "for" and "while" so that FXC does not
attempt to unroll them. This is to work around an FXC bug where it fails
to unroll loops with gradient operations.
FXC ostensibly unrolls such loops because gradient operations require
uniform control flow, and loops that have varying iterations may
possibly not be uniform. Tint will eventually validate that control flow
is indeed uniform, so forcing FXC to avoid unrolling in these cases
should be fine.
Bug: tint:1112
Change-Id: I10077f8b62fbbb230a0003f3864c75a8fe0e1d18
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69880
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
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>
Spread the array zeroing across as many workgroup invocations as possible.
Bug: tint:910
Change-Id: I1cb5a6aaafd2a0a4093ea3b9797c173378bc5605
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60203
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@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>
If the array size is greater than a threshold.
This is a work around for FXC stalling when initializing large arrays
with a single zero-init assignment.
Bug: tint:936
Fixed: tint:943
Fixed: tint:942
Change-Id: Ie93c8f373874b8d6d020d041fa48b38fb1352f71
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56775
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Antonio Maiorano
James Price
Ben Clayton