James Price
5c61d6d12c
msl: Fold &* when converting module-scope vars
...
This transform was previously converting this code:
```
var<private> v : f32;
fn foo() {
bar(&v);
}
```
into this:
```
fn foo(vp : ptr<private, f32>) {
bar(&*vp); // Invalid, since ptr args must be &ident
}
```
Fixed: tint:1086
Change-Id: Ic9efafa219c89a11a4d6e1d11fc69b3c0b9a5464
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60520
Kokoro: Kokoro <noreply+kokoro@google.com >
Auto-Submit: James Price <jrprice@google.com >
Reviewed-by: Ben Clayton <bclayton@google.com >
Commit-Queue: Ben Clayton <bclayton@google.com >
2021-08-04 19:18:38 +00:00
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
