Commit Graph

23 Commits

Author SHA1 Message Date
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 f24b37e122 writer/msl: Rework string printing
Add `out` parameters to expression and type generators.

Use the new helper classes in TextGenerator.
Cleans up bad formatting.

Prepares the writer generating 'pre' statements, required for atomics.

If-else statements are generated slightly differently. This is done so that 'pre' statements for the else conditions are scoped correctly. This is identical to the HLSL writer.

Bug tint:892

Change-Id: I4c6e96c90673ba30898b3682bf3198497d63a2d4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56067
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
2021-06-28 15:30:57 +00:00
Ben Clayton 41f21fe05b writer/hlsl: Emit zero values for private variables
Also remove the unreachanble constructor logic in EmitHandleVariable.
Variables of the handle storage class cannot have initializers.

Fixed: tint:173
Change-Id: I7c997a8b6a70308ff9b5c42fa1198810ee365bac
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55258
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
2021-06-21 19:37:58 +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 165512c57e writer/hlsl: Emit UBO as an array of vector
Instead of a ConstantBuffer.

HLSL requires that each structure field in a UBO is 16 byte aligned.
WGSL has much looser constraints with its UBO field alignment rules.

Instead generate an array of uint4 vectors, and index into this, much
like we index into [RW]ByteAddressBuffers for SSBOs.

Extend the DecomposeStorageAccess transform to support uniforms too.
This has been renamed to DecomposeMemoryAccess.

Change-Id: I3868ff80af1ab3b3dddfbf5b969724cb87ef0744
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55246
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-06-18 21:15:25 +00:00
James Price 567f2e4f3b transform/msl: Run InlinePointerLets and Simplify
This will be relied on by the upcoming arrayLength transform.

Update test expectations.

Change-Id: Ib74b647abcd6f4393f9899ce40bbf06f6e53e7f4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55180
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-06-18 09:47:23 +00:00
David Neto 17287fcf1a spirv-reader: Set workgroup size, but not specializable
The WorkgroupSize builtin decoration applies to a composite constant.
Because WGSL does not yet support specializable constants for this,
use the *default* values for that SPIR-V spec constant.

Update end-to-end test expectations.

Fixed: tint:503
Change-Id: I012b316d13544ab9282e3276b58906327adab133
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/41960
Auto-Submit: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: Alan Baker <alanbaker@google.com>
2021-06-17 22:40:43 +00:00
David Neto 1e19b55d19 spirv-reader: switch to HLSL-style pipeline IO
- When storing to sample_mask output, write to the 0th element
- Only make a return struct if it has members
- Adjust type signedness coercion when loading special builtins.
- Adapt tests

- Update expectations for end-to-end tests

- Handle sample_mask with stride
  Input variables normally don't have layout. But they can have it
  up through SPIR-V 1.4.
  Handle this case in the SPIR-V reader, by seeing through the
  intermediate alias type created for the strided array type.

Bug: tint:508
Change-Id: I0f19dc1305d3f250dbbc0698a602288c34245274
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54743
Auto-Submit: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
2021-06-17 09:10:04 +00:00
Ben Clayton d47eb3a965 writer/hlsl: Generate padding for UBO padded structs
Combined with the new PadArrayElements transform, arrays with strides
are now correctly emitted.

Fixed: tint:182
Fixed: tint:895
Change-Id: I26a1be94dee6e4c9d9747c8317a932fc1fb3c810
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54640
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-06-16 09:50:11 +00:00
Ben Clayton 5d2f34ecf2 writer/hlsl: Simplify emission logic, clean up output
And fix issues where global variables would not be emitted unless they were transitively referenced by an entry point.

This change requires crbug.com/tint/697 to be fixed before landing.

Change-Id: I712bd9d369e08c9a3cdfb0f114c3609584f91f28
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54241
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-06-16 09:19:36 +00:00
James Price 2940c7002c writer/msl: Avoid generating unnecessary pointers
When moving private and workgroup variables into the entry point,
generate pointers to pass as arguments to sub-functions on demand,
instead of upfront. This removes a bunch of unnecessary dereferences
for accesses inside the entry point, and one function variable.

Change-Id: I7d1aabdf14eae33b569b3316dfc0f9fbd288131e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54300
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
2021-06-11 12:29:56 +00:00
Ben Clayton 0aa7edbbd5 transform/spirv: Use InlinePointerLets & Simplify
arrayLength() will take a pointer to a storage buffer array.
This pointer may pass through function scoped let statements.
To make this intrinsic easier to generate, inline the pointer lets and
remove chains of &*&*.

Bug: tint:806
Change-Id: Ib2c79a9c38cba7391cbb4313986af9a72b0f0435
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54060
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
2021-06-10 17:34:44 +00:00
Ben Clayton 1858854f7e Add optional access to ptr<>
This also completes the work to resolve the access controls for each
storage type.

Fixed: tint:846
Change-Id: Iab24057ec14620a2978ec63c4a91ba12d1bc6e9b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53381
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-06-04 22:17:37 +00:00
Ben Clayton 93e8f527ee wgsl: Deprecate [[access]] decorations
Handle access control on var declarations instead of via [[access]]
decorations. This change does the minimal work to migrate the WGSL
parser over to the new syntax. Additional changes will be needed
to correctly generate defaulted access qualifiers, as well as
validating access usage.

The [[access]] decorations are still supported by the WGSL parser,
with new deprecated warnings, but not for aliases. Example:
   var x : [[access(x)]] alias_to_struct;

Making this work is far more effort than I want to dedicate to backwards
compatibility, and I do not beleive any real-world usage will be doing
this.

Still TODO:
* Adding access control as the optional, third parameter to ptr<>.
* Calculating default accesses for the various storage types.
* Validating usage of variables against the different accesses.

Bug: tint:846
Change-Id: If8ca82e5d16ec319ecd01f9a2cafffd930963bde
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53088
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
2021-06-04 20:41:47 +00:00
James Price 5102f87e38 writer/msl: Add parentheses for member accesses
The LHS should be wrapped in parentheses if it has lower precedence
than the access. This fixes issues with pointer dereferences followed
by member accesses, where we were previously generating *a.b.

Fixed: tint:831
Change-Id: I8a194ad4f54c80a01c24eb983ec8064037575216
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51963
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-05-27 14:15:47 +00:00
James Price 7a47fa8495 writer/msl: Handle private and workgroup variables
Add a transform that pushes these into the entry point and then passes
them by pointer to any functions that need them.

Since WGSL does not allow non-function storage class at
function-scope, add a DisableValidation attribute to bypass this
check.

Fixed: tint/726
Change-Id: Ic1f4cd691a54c19e77a60e8ba178508e4249bfd9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51962
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-05-26 15:41:02 +00:00
James Price e5fdd58352 writer/msl: Emit const on array and ptr parameters
This fixes issues with passing constant arrays to functions.

Change-Id: I6e2f1c3f64df836c0b6a55ab925cf3c2bc317733
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51861
Auto-Submit: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-05-22 12:40:04 +00:00
Ben Clayton ed86bf99b0 Add transform::Simplify
Performs basic peephole optimizations on the AST.

Use in transform::Hlsl.

Required to have the DecomposeStorageAccess transform operate correctly with the output of InlinePointerLets transform, specifically when declaring `let` pointer expressions to storage buffers.

Fixed: tint:221
Fixed: tint:492
Fixed: tint:829
Change-Id: I536390921a6492378104e9c3c100d9e761294a27
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51921
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
2021-05-21 21:01:23 +00:00
Ben Clayton ada560b289 writer/hlsl: Fix continuing block emission
Inline the `continuing` block in the places where `continue` is called.

Simplifies the emission, and fixes emission of `let` statements in the loop.

Also fix random indenting of intrinsic functions.

Fixed: tint:744
Fixed: tint:818
Change-Id: I06994dbc724bc646e0435a1035b00760eaf5f5ab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51784
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
2021-05-20 21:31:37 +00:00
Ben Clayton 06aa88aa97 hlsl: Pointer support
Add `transform::InlinePointerLets` - a Transform that moves all usage of function-scope  `let` statements of a pointer type into their places of usage.

Make the HLSL writer transform pointer parameters to `inout`.

Fixed: tint:183
Change-Id: I0a7552fa6cd31c7b7691e64feae3170a81cc6c49
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51281
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-05-19 19:16:32 +00:00
James Price cbe816f93d writer/msl: Generate address spaces for pointers
Add more E2E tests to cover pointers with different storage classes.

Fixed: tint:815
Change-Id: I224a794cdf60648ce71dc9a0922d489542995be1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51404
Auto-Submit: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
2021-05-19 10:38:18 +00:00
Ben Clayton 09d53d5ed1 writer/hlsl: Emit `static` for private global variables
Fixed: tint:812
Change-Id: I822eb6b13f07a74eb47c4c8e18d1f75dcc4818bf
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51366
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-05-19 09:03:48 +00:00
Ben Clayton 9b54a2e53c Implement Pointers and References
This change implements pointers and references as described by the WGSL
specification change in https://github.com/gpuweb/gpuweb/pull/1569.

reader/spirv:
* Now emits address-of `&expr` and indirection `*expr` operators as
  needed.
* As an identifier may now resolve to a pointer or reference type
  depending on whether the declaration is a `var`, `let` or
  parameter, `Function::identifier_values_` has been changed from
  an ID set to an ID -> Type* map.

resolver:
* Now correctly resolves all expressions to either a value type,
  reference type or pointer type.
* Validates pointer / reference rules on assignment, `var` and `let`
  construction, and usage.
* Handles the address-of and indirection operators.
* No longer does any implicit loads of pointer types.
* Storage class validation is still TODO (crbug.com/tint/809)

writer/spirv:
* Correctly handles variables and expressions of pointer and
  reference types, emitting OpLoads where necessary.

test:
* Lots of new test cases

Fixed: tint:727
Change-Id: I77d3281590e35e5a3122f5b74cdeb71a6fe51f74
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/50740
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-05-18 10:28:48 +00:00