Commit Graph

10 Commits

Author SHA1 Message Date
Ben Clayton 89a0bde59c transform: Optimize ZeroInitWorkgroupMemory for arrays
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>
2021-07-30 14:08:06 +00:00
Ben Clayton e027e81bf2 writer/hlsl: Emit helper functions for storage class atomic intrinsics
By generating a helper function for these, we can keep the atomic expression pre-statement-free. This can help prevent for-loops from being transformed into while loops, which can upset FXC.

We can't do the same for workgroup storage atomics, as the InterlockedXXX() methods have the workgroup-storage expression as the first argument, and I'm not aware of any way to make a user-declared parameter be `groupshared`.

Change-Id: I8669127a58dc9cae95ce316523029064b5c9b5fa
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57462
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-07-09 16:50:14 +00:00
Ben Clayton 1843c0b8d7 [writer/hlsl]: Fix order of atomic method arguments.
Change-Id: Ice1b07c748bc6502a51b29690dfc00466a684c12
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57461
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-07-09 12:35:54 +00:00
Sarah e6cb51e715 validation: compute shader must include 'workgroup_size' in its attributes
Bug: tint:884
Change-Id: If96c6df3247fee142a779117fa26d006afd4f7ef
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55680
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-06-29 18:39:44 +00:00
Ben Clayton f2ec7f38e5 writer/msl: Implement atomics
Common logic between the HLSL, WGSL and MSL writers has been moved into
the TextGenerator base class.

Fixed: tint:892
Change-Id: I0f469516947fe64817ce6251e436da74e5e176e8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56068
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
2021-06-29 11:53:15 +00:00
Ben Clayton 75db82c96b sanitizers: Use the ZeroInitWorkgroupMemory transform
Zero the workgroup memory for all backends.
We can probably disable this for the backends that support workgroup zeroing, but that's an optimization we can perform later.

Fixed: tint:280
Change-Id: I9cad919ba3a15b8cedfe6939317d1f6b95425453
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55244
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-06-18 22:44:31 +00:00
Ben Clayton c3dc300fcb writer/spirv: Implement atomics
Bug: tint:892
Change-Id: Ic0de538c76fd7cfe8fd3d7c25d2d61dd74aa1494
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54658
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-06-18 21:15:25 +00:00
Ben Clayton e6d171ac66 writer/hlsl: Implement atomics
Storage buffers are emitted as `ByteAddressBuffer`s in HLSL, so we have to jump through hoops to support atomic ops on storage buffer atomics.
Workgroup atomics are far more conventional, but very little code can be shared between these two code paths.

Bug: tint:892
Change-Id: If10ea866e3b67a093e87aca689d34065fd49b705
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54651
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-06-18 18:56:13 +00:00
Ben Clayton 0a32a724f4 writer/wgsl: Emit atomic types
Bug: tint:892
Change-Id: Ie483167bcf669e5f2d6b5489a915584fc3678183
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54649
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-06-18 18:56:13 +00:00
Ben Clayton 6a77236d8c intrinsics.def: Add atomic intrinsics
Change-Id: I26a9ce9e1978aa2542ce2b3f17c9f5861e556a8a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54657
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-06-18 18:56:13 +00:00