Commit Graph

13 Commits

Author SHA1 Message Date
Antonio Maiorano e5dbe24e94 MSL writer: make signed int overflow defined behaviour
Bug: tint:124
Change-Id: Icf545b633d6390ceb7f639e80111390005e311a1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60100
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: David Neto <dneto@google.com>
2021-07-29 13:51:47 +00:00
Ben Clayton 0e41747281 writer/hlsl: Emit for-loops where possible.
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>
2021-07-02 22:17:25 +00:00
Ben Clayton 2bb45389b7 writer/hlsl: Zero initialize with (T) 0
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>
2021-07-02 19:27:42 +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 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
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 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
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
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