These work on scalar and vector of bool, and map to ast::BinaryOp::kOr
and kAnd.
Bug: tint:1043
Change-Id: I009edf8e43c21cb75ccfdcde1106ec177d2fe50e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59561
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: David Neto <dneto@google.com>
Add end to end tests
Fixed: tint:1026
Change-Id: I10813cbe6dc4f1bccddf9a8a29e3a249a364c051
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59663
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Also move:
test/fxc_bugs/vector_assignment_in_loop
to
test/bug/fxc/vector_assignment_in_loop
Change-Id: I7bbfc476fdb7a3296025609625e322fed8d16285
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59444
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>
Implement these for all the writers.
SPIR-V reader not implemented (the old overloads weren't implemented either).
Deprecate the old overloads.
Fixed: tint:54
Change-Id: If66d26dbac3389ff604734f31b426abe47868b91
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59302
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
This adds SPIR-V assembly and WGSL tests derived from VK-GL-CTS commit
571256871c2e2f03995373e1e4a02958d8cd8cf5. The following procedure was
followed:
- Those .amber files in VK-GL-CTS wholly owned by Google were
identified
- All GLSL and SPIR-V shaders were extracted from the Amber files and
converted into SPIR-V binaries
- The compact-ids pass of spirv-opt was applied to each binary
- Duplicate binaries were removed
- spirv-opt -O was used to obtain an optimized version of each remaining
binary, with duplicates discarded
- Binaries that failed validation using spirv-val with target
environment SPIR-V 1.3 were discarded
- Those binaries that tint could not successfully convert into WGSL were
put aside for further investigation
- SPIR-V assembly versions of the remaining binaries are included in
this CL
- test-runner with -generate-expected and -generate-skip was used to
generate expected .spvasm, .msl, .hlsl and .wgsl outputs for these
SPIR-V assembly tests
- Each successfully-generated .expected.wgsl is included in this CL
again, as a WGLSL test
- test-runner with -generate-expected and -generate-skip was used again,
to generate expected outputs for these WGSL tests
Change-Id: Ibe9baf2729cf97e0b633db9a426f53362a5de540
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58842
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
https://github.com/gpuweb/gpuweb/pull/1945 changes the SPIR-V mapping of this operator so that it now maps to OpFRem instead of OpFMod. Polyfill OpFMod with `x - y * floor(x / y)`
Also map the MSL output of this operator to use `fmod()`.
Behavior of this operator is now consistent across all backends.
Fixed: tint:945
Fixed: tint:977
Fixed: tint:1010
Change-Id: Iefa009b905989c55ace24e073ab0e261c7cf69b0
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58393
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Change-Id: Ied931db182c11f26bffda4900641e41a1e9b5ee8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59027
Commit-Queue: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
A hypothetical case discussed in the WG call, which I wasn't entirely sure was going to work. It does.
Change-Id: I855f1e02af2fe62214e3fb964f6937e3d88016eb
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59021
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
The only non-trivial case is SPIR-V, which generates OpSelect to
choose between 1 or 0.
Fixed: tint:997
Change-Id: Ifda7b3ec1e0a713843a2da7ed59c3449d4eec8bd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58521
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
The TypeConstructorExpression logic that tested for splats was not considering references. This led to broken emission for the SPIR-V and HLSL backends.
Fixed: tint:992
Change-Id: I9824b71f526997f91d380c09b459f4fd73065b19
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58397
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Tests should be as fine-grained as possible.
Not all the permutations were covered.
Not all the backends were actually generating anything useful, as the
functions were not called by an entry point.
Change-Id: I76658533d3112d202f7482a19531079550e66f83
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58392
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Arrays can be extremely large, and having the load and store functions unroll the elements can make the complier explode.
Fixed: chromium:1229233
Change-Id: Ieb5654254e16f5ce724a205d21d954ef9a0cd053
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58382
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
For loop initializers and continuing statements do not have a BlockStatement as their parent.
Handle removal of these statements with a new Transform::RemoveStatement() helper
Fixed: tint:990
Change-Id: I24e7b18dcf71d3ef0a4d3ee68b9f68518e0eb5e8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58063
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Much like sem::Type, it greatly simplifies downstream logic if we can compare sem::Intrinsic pointers to know if they refer to the same intrinsic overload.
Change-Id: If236247cd3979bbde821d9294f304ab85ba4938e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58061
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
And call these helpers instead of inlining complex statements.
Cleans up output, and helps prevent for-loops decaying to while loops.
Change-Id: I6ac31b18ce6c5fac0e54e982f7db3bb298f7edb2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58060
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
CloneContext::Replace(T* what, T* with) is bug-prone, as complex transforms may want to clone `what` multiple times, or not at all. In both cases, this will likely result in an ICE as either the replacement will be reachable multiple times, or not at all.
The CTS test: webgpu:shader,execution,robust_access:linear_memory:storageClass="storage";storageMode="read_write";access="read";atomic=true;baseType="i32"
Was triggering this brokenness with DecomposeMemoryAccess's use of CloneContext::Replace(T*, T*).
Switch the usage of CloneContext::Replace(T*, T*) to the new function form.
As std::function is copyable, it cannot hold a captured std::unique_ptr.
This prevented the Replace() lambdas from capturing the necessary `BufferAccess` data, as this held a `std::unique_ptr<Offset>`.
To fix this, use a `BlockAllocator` for Offsets, and use raw pointers instead.
Because the function passed to Replace() is called just before the node is cloned, insertion of new functions will occur just before the currently evaluated module-scope entity.
This allows us to remove the "insert_after" arguments to LoadFunc(), StoreFunc(), and AtomicFunc().
We can also kill the icky InsertGlobal() and TypeDeclOf() helpers.
Bug: tint:993
Change-Id: I60972bc13a2fa819a163ee2671f61e82d0e68d2a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58222
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This is a reland of fd5829e5ea
Original change's description:
> Validate storage class constraints
>
> As defined by https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints
>
> Bug: tint:643
> Change-Id: I9c78ba69a792a80c263a17b0a6e9b4810fdb7f30
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56780
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
> Reviewed-by: Ben Clayton <bclayton@google.com>
Bug: tint:643
Change-Id: I62036d615a062597339a9d130b7ccf49b5be26c7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58120
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This uses FXC compilation failure mitigation for _any_ vector index assignment that has a non-constant index. FXC can still fall over if the loop calls a function that performs the dynamic index.
Use some vector swizzle logic to avoid branches in the helper.
Fixed: tint:980
Change-Id: I2a759d88a7d884bc61b4631cf57feb4acc8178de
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57882
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
The UBO must have a stride that is a multiple of 16 bytes.
Note that this change was part of https://dawn-review.googlesource.com/c/tint/+/56780
but the CL was reverted because it broke Dawn. This CL relands part of
the change, and adds the macro TINT_EXPECTS_UBOS_TO_BE_MULTIPLE_OF_16 so
that Dawn can conditionally compile against it.
Bug: tint:984
Bug: tint:643
Change-Id: I303b3fe81ff97c4933c489736d5d5432a59ce9b7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57921
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This change introduces sem::CompoundStatement, a new base class for
statements that can hold other statements.
sem::BlockStatements now derives from sem::CompoundStatement, and
this change introduces the following new CompoundStatements:
* `sem::IfStatement`
* `sem::ElseStatement`
* `sem::ForLoopStatement`
* `sem::LoopStatement`
* `sem::SwitchStatement`.
These new CompoundStatements are now inserted into the semantic
tree as now documented in `docs/compound_statements.md`.
The `sem::BlockStatement::FindFirstParent()` methods have been
moved down to `sem::Statement`.
The `Resolver::BlockScope()` method has been replaced with
`Resolver::Scope()` which now maintains the `current_statement_`,
`current_compound_statement_ ` and `current_block_`. This
simplifies statement nesting.
The most significant change in behavior is that statements now
always have a parent, so calling Block() on the initializer or
continuing of a for-loop statement will now return the
BlockStatement that holds the for-loop. Before this would
return nullptr.
Fixed: tint:979
Change-Id: I90e38fd719da2a281ed9210e975ab96171cb6842
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57707
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
The original problem appears to be fixed.
Fixed: tint:219
Change-Id: I8d16fbb715da3ca149769699c86f86a4bed85b4f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57703
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This class is very large, so factoring it out into its own seperate
file.
Also renaming it to InspectorBuilder, in anticipation of adding tests
that don't build their state through explicitly calling into the
ProgramBuilder.
BUG=tint:967
Change-Id: I316458a969479200edf6962ab2008180f9a1e7e2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57800
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Ben Clayton <bclayton@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>
Move the bulk of the constant evaulation logic out of transform::FoldConstants and into Resolver and sem::Expression.
transform::FoldConstants now replace TypeConstructor nodes that have a constant value on the expression.
This is ground work to:
* Cleaning up the HLSL uniform buffer indexing, which is `/` and `%` arithmatic heavy
* Prepares us to handle `constexpr` when it lands in the spec
* Provide a centralized place to do constant evaluation, instead of the
having similar logic scattered around the codebase.
Change-Id: I3e2f542be692046a8d243b62a82556db519953e7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57426
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
This reverts commit fd5829e5ea.
Reason for revert: Temporarily reverting as this is preventing a tint->dawn roll, which is needed to fix the dawn->chrome roll.
Original change's description:
> Validate storage class constraints
>
> As defined by https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints
>
> Bug: tint:643
> Change-Id: I9c78ba69a792a80c263a17b0a6e9b4810fdb7f30
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56780
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
> Reviewed-by: Ben Clayton <bclayton@google.com>
TBR=bclayton@google.com,amaiorano@google.com,noreply+kokoro@google.com,tint-scoped@luci-project-accounts.iam.gserviceaccount.com
Change-Id: I8dbd0e46b3e3291ef08797a196d0d9abd2a78845
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:643
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57704
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Report whether one was generated so that Dawn knows to use the
`-fpreserve-invariance` compiler option.
Bug: tint:772
Change-Id: Ife1eb05265646727dc864f12f983781af4df3777
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57644
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Make sure the other backends ICE on unrecognized attributes.
Add E2E tests, currently skipped for the other backends.
Bug: tint:772
Change-Id: I4e68d111ff79b19ebb6c574058a91debcb746011
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57642
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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.
Change-Id: Id034ea5ea9be601661ddb78db973015d845c420f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57463
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
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>
And remove the u32 overload of frexp (it's not in the spec).
Brings the number of failing tint end to end tests for MSL down to 19/1098.
The WG still haven't found consensus on reworking these two intrinsics.
It's very likely that their signature will change so that they return a structure instead of returning a value and outputing another as a pointer.
Until the WG makes a decision, let's implement these according to the current spec.
Some overloads are still failing due to MSL missing overloads of the pointer parameter being in the `threadgroup` address space.
I'm holding off fixing these until we know what's happening with these intrinsics.
See also:
https://github.com/gpuweb/gpuweb/issues/1480https://github.com/gpuweb/gpuweb/issues/1846
Change-Id: Ib6764e6659d840db41bc65fed2b8b283d1056c3d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57421
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
The level parameter needs to be zero for a number of texture overloads.
Bodging the template to emit 0 instead of 1 for any `level : i32` parameter is the easiest fix here.
Change-Id: I69222578efcd0d4f4f267fb59fec691b52786d9c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57301
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Removes the unneeded texture_external overload of textureSample from
intrinsics.def.
Bug: tint:858
Change-Id: I49935bae47542af7b440c74c96418e005daf9e19
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53940
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
This reverts commit 26b6edc545.
Reason for revert: Seems to be breaking Dawn's tests. Investigation required.
Original change's description:
> writer/hlsl: Special case negative zero
>
> Fixed: tint:960
> Change-Id: I060bc6b7a9ad4d21dd5cadb4b68998c7e54ebaed
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57142
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Commit-Queue: Ben Clayton <bclayton@google.com>
> Auto-Submit: Ben Clayton <bclayton@google.com>
> Reviewed-by: James Price <jrprice@google.com>
# Not skipping CQ checks because original CL landed > 1 day ago.
Change-Id: Ia0b0ec996f2ed6b1599a344c970f155c12314ea9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57460
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
These were removed from the spec in:
https://github.com/gpuweb/gpuweb/pull/1914
Bug: tint:921
Change-Id: I4e584fdee9cf540a192f12d1208595056e081410
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57300
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
Transforms a for-loop into a loop.
Will be required by the SPIR-V writer.
Bug: tint:952
Change-Id: Iba859bd144d702cee85374f2cfcb94cd7465ca98
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57202
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Runs the statement(s) at the end of the lexical scope
Change-Id: I74de02b7204b56016c7bbe71fee4ece498d3570d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51923
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
WGSL:
* Remove vertex_idx and instance_idx.
These are now vertex_index and instance_index.
It seems this was removed once before, then reverted due to CTS
failures, but the original change never landed again.
* Remove the [[set(n)]] decoration. This has been [[group(n)]] for
months now.
API:
* Remove deprecated enums from transform::VertexFormat.
* Remove transform::Renamer constructor that takes a Config. This should
be passed by DataMap.
* Remove ast::AccessControl alias to ast::Access.
Change-Id: I988c96c4269b02a5d77163409f261fd5923188e0
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56541
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Also add missing msl macros to the renamer.
Bug: tint:951
Change-Id: I543e6eae885c979596ca63f9d6c7378dd5425e8a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56941
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
And add `vec` and `mat` to the reserved keyword list (see https://github.com/gpuweb/gpuweb/pull/1896)
Move these reserved keyword checks out of the lexer and into the parser.
Generate a sensible error message.
Add tests.
Change-Id: I1876448201a2fd773f38f337b4e49bc61a7642f7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56545
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@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>
We now generate valid code for this.
Change-Id: I9402c6aa32365fca5683e91b3fa345ed7d3e34ed
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56622
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
FXC has trouble dealing with these.
This was originally added to handle returning arrays as structures.
HLSL supports typedefs, which is a much simpiler solution, and doesn't upset FXC.
Bug: tint:848
Bug: tint:904
Change-Id: Ie841c9c454461a885a35c41476fd4d05d3f34cbf
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56774
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
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>
A Transform that attempts to convert WGSL `loop {}` statements into a for-loop statement.
For-loops cause less broken behavior with FXC than our current loop constructs.
Bug: tint:952
Change-Id: I0b7b1dbec9df4c004b0419d3feae53a195ec79bb
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56767
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
This transform is intended to clean up the output of the SPIR-V reader, so that we can pattern match loops that can be transformed into a for-loop.
Bug: tint:952
Change-Id: Iba58e4e1f6e20daaf7715e493df53346cdb7c89f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56766
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
These indices were a mix of signed and unsigned.
Modulus on the signed integers was producing FXC warnings about performance.
Change-Id: Ib82f4296199a09d2f03be8b06314feefce0022e2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56765
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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>
Fixed many tests that had empty structures.
Change-Id: Id91312afa39a6293426f99d0dd12578dba46aa61
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56621
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>
UniqueIdentifier() will generate a program-global unique symbol.
MslGeneratorImplTest.AttemptTintPadSymbolCollision tests for collisions with the field names.
TextGeneratorTest.UniqueIdentifier_ConflictWithExisting tests for collisions between general symbols.
Fixed: tint:654
Change-Id: If2ba75d04ff0e2a9975e878596ac114d51adcd46
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56580
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Sarah Mashayekhi <sarahmashay@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
These test primarily test emission of the variables based on transitive function call usage.
Change-Id: I0f64cb6496ed1238000cb8a6c97a1a445de0ab41
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56546
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
We will want this transform to do more bounds and argument sanitization.
Bug: tint:748
Change-Id: I38cb9623622e9f5ab85d8cd420d669ca6be77099
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56543
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
When indexing into vectors in a loop, FXC sometimes fails to determine
the max number of iterations when attempting to unroll the loop,
resulting in "error X3511: forced to unroll loop, but unrolling
failed.". We work around this by calling a function that sets the input
value at the input index into an inout vector. This seems to nudge FXC
enough for it to determine the number of loop iterations to unroll.
Bug: tint:534
Change-Id: I52cb209be29fcad8fbb91283c7be8c6e22e00656
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56140
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Use the sanitizer to add the decoration only to the variables that are
vertex outputs and fragment inputs.
Bug: dawn:963, tint:746
Change-Id: I1b91cf3550fb3c6f583d69e822444534a576e0cd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56460
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Vulkan requires that shader inputs/outputs that are integers must be
decorated with Flat.
Bug: tint:746, dawn:956
Change-Id: I648451b9aa559d08415bada904dee5f9d1e4e60f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56400
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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>
WGSL supports select() with vectors, where the condition is a
scalar. To support this in SPIR-V versions older than 1.4, we need to
splat the condition operand to a vector of the same size as the
objects.
Fixed: tint:933
Change-Id: I571af46e74cd7bb24093524ccfed25a3ed612676
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56340
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>
Add the SampleRateShading capability if the sampling type is `sample`.
Bug: tint:746
Change-Id: I20fb25913f5c0919b6d16a9d0e9fc8b1551ff7ea
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56247
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Add E2E tests to cover all of the parameter combinations.
Mark the attribute as unimplemented in the other backends.
Bug: tint:746
Change-Id: I86881ff0b224fe93670db42473341ae185eeabdd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56244
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
These are expressions, and not specific to `var`s.
Break the tests up into finer granularity.
Bug: tint:656
Change-Id: I6873407127871dfaec55d90f02e0490eef929a30
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56071
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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>
Remove `pre` and `out` parameters from most generator methods.
Use the `out_` string stream in TextGenerator, add helpers to TextGenerator to simplify line printing.
Remove the `pre` and `out` fields from TestHelper.
Cleans up the `pre` aspects of the HLSL writer, so the same concept can be used by the MSL writer.
Fixes indentation bugs in formatting.
Change-Id: Ia35daf632c7c7b84a6fbf3b9ae42baaeb3c97649
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55960
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
HLSL usually implicitly casts a vector down to a scalar, but this breaks when passing the vector to RWByteAddressBuffer.Store (for DXC only).
Fixed: tint:827
Change-Id: I67d0bc6e9185de3d434a7aaeb575d83850111ec5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55760
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This fixes the SPIR-V and MSL tests for these intrinsics.
Change-Id: Id6f48682285ff17cb1fa7ef618f34b02f553332b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55681
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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>
Varaibles can infer types now, in which case the type_ field is null.
Fixed: chromium:1221120
Change-Id: I0cb2a6a2e8128c56625f48940cf73cf4cadb22ce
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55252
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Inline the `continuing` block in the places where `continue` is called.
Simplifies the emission, and fixes emission of `let` statements in the loop.
This fix matches the same approach in writer/hlsl.
See: https://dawn-review.googlesource.com/c/tint/+/51784
Fixed: tint:833
Fixed: tint:914
Change-Id: If4d8cde62dfaf8efa24272854ca7ff5edc0a8234
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55341
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Generate a uniform buffer that will receive the lengths of all storage
buffers, and use this to implement calls to arrayLength(). The
transform is provided with a set of mappings from storage buffer
binding points to the corresponding index into the array of buffer
lengths. The transform reports whether it generated the uniform
buffers or not.
Use this transform from the MSL sanitizer, using the binding number as
the index into the array. This matches the behavior of spirv-cross,
and so works with how Dawn already produces this uniform buffer.
Bug: tint:256
Change-Id: I2682d2d024e8daa30f78270b8cfb6bbb32632133
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54480
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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>
Zero initializes all referenced workgroup storage classed variables used by each entry point.
Bug: tint:280
Fixed: tint:911
Change-Id: I3fca26a10f015f08fedef404720bbe6fd7b343a9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55243
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
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>
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>
Instead of just generating pointers to functions, generate pointers to all permuted storage types and accesses.
Change-Id: I930b20150d823877eaf72dd7cac850078fe22f35
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54656
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: David Neto <dneto@google.com>
After implementing validation and fairly exhaustive tests, discovered
that conversion of scalar vector to bool vector did not work in the
spir-v backend. For module scope variables, we use and rely on the
FoldConstants transform to ensure no conversion needs to take place.
This is necessary because we cannot easily introduce temporary values
and refer to them when casting at module scope. Note that for the same
reason, module-level conversions are always constant foldable, so this
works. For function-level conversions, implemented support to emit a
comparison against a zero value, and store the result in the bool
vector.
Bug: tint:865
Change-Id: I0528045e803f176e03428bc7eac31ae06920bbd7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54744
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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>
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>
- 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>
This translates to/from OpNot for SPIR-V, and ~ for all three textual
language backends.
Fixed: tint:866
Change-Id: Id934fb309221e3fca0e7efa33edaaae137fd8085
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54980
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>
Dynamic indexes are limited to references to matrices and arrays
https://github.com/gpuweb/gpuweb/pull/1801
Bug: tint:867
Change-Id: I114daa053c8a4ffd23ce784ac4538567a551cc21
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54701
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Sarah Mashayekhi <sarahmashay@google.com>
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>
Replaces arrays with an explicit stride with an array to a structure holding the element padded with a `[[size]]` decoration.
Note that the HLSL writer is still not correctly emitting structure fields with a `[[size]]`, which will be fixed in a follow up change.
Bug: tint:182
Bug: tint:895
Fixed: tint:180
Fixed: tint:649
Change-Id: Ic135dfc89309ac805507e9f39392577c7f82d154
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54582
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
These have been deprecated, and their usages in Dawn, CTS and samples have been updated.
Fixed: tint:846
Change-Id: I74b831fd5be2e7ca02e8208835eac8beddcef9af
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54325
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Fixes issues with using arrays as function return types.
Fixed: tint:848
Change-Id: Iee8af0f2cea9d19e448176446c6599be2bd32316
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54321
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
And replace the MSL writer's logic to do this with the transform.
We need to do the same thing in HLSL, and in the future GLSL too.
Partially reverts fbfde720
Change-Id: Ie280e011bc3ded8e15ccacc0aeb12da3c2407389
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54242
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
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>
We have the end-to-end test-runner which validates all this stuff.
There's no need to also Validate in the unit tests.
Change-Id: I8fbc1ecd395efcedec4aa41085e691c88f3b66a5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54160
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
Also split out validation tests from call_test.cc into call_validation_test.cc.
Bug: tint:886
Change-Id: I1e1dee9b7c348363e89080cdecd3119cc004658f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54063
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Wrap the texture expression in parentheses when it has lower
precendence than the function call operator.
Cast integer coordinates to unsigned integers as required by MSL.
Fixed: tint:536
Change-Id: I957e6be3c51044959e25e0be96c2d2c65db18187
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53962
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Move these module-scope variables to entry point parameters and pass
them as arguments to functions that use them. Disable entry point IO
validation for them.
Emit [[texture()]] and [[sampler()]] attributes on these entry point
parameters.
Fixed: tint:145
Change-Id: I936a80801875a5d0b6cd98a2e8f3e297a2f53509
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53961
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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>
Some now pass, some still fail.
If we fix bugs and they begin to pass, we have a better chance of noticing and marking the issues as fixed.
Change-Id: I7278f960967a570ccde8865dcf2aa580235559bf
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54062
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
The old non-pointer argument overload remains, but is now deprecated.
Bug: tint:806
Change-Id: Ic917ccec0f162414ce1b03df49e332ad84d8060f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54061
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
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>
- Add resolver/call_test.cc for new unit tests, and move a couple that
were in resolver/validation_test.cc to it
- Fix CalculateArrayLength transform so that it passes the address of
the u32 it creates to the internal function
- Fix tests broken as a result of this change
Bug: tint:664
Change-Id: If713f9828790cd51224d2392d42c01c0057cb652
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53920
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
The argument order between WGSL and SPIR-V is different (condition is first in SPIR-V, last in WGSL)
Fixed: tint:560
Change-Id: I56c659c441292e05f71a24d96dbc9f93f25b71f0
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53620
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Much like uintN, intN, floatN - boolN is far more common, and easier to read than vector<bool, N>
Change-Id: I51f9edc003c590266316d3eba286ca2f6882da10
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53390
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
This does not pass validation.
Change-Id: I3e63ad899d346321a4dea13f3a681d01546ed2bb
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53621
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
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>
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>
This is required to generate valid MSL code, and will soon be
validated by Tint too.
Change-Id: I4c5f5c4ecb1c91131c934de1132217d9f6be1f8e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53420
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Add a config parameter for the CanonicalizeEntryPoint transform that
selects between emitting builtins as parameters (for MSL) or struct
members (for HLSL).
This fixes all of the shader IO issues in Tint's E2E tests for MSL.
Fixed: tint:817
Change-Id: Ieb31cdbd2e4d96ac41f8d8515fd07ead8241d770
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53282
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
This fixes constructors for structures that contain padding members
due to explicit layout attributes.
Also fix one test that was wrongly using an identity type constructor
for a structure.
Fixed: tint:853
Change-Id: I0a3e84fcd7c6a7f2ad92a4970ed11378e6ce2465
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53240
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Use the new [[stage()]] decorations in intrinsics.def to validate that intrinsics are only called from the correct pipeline stages.
Fixed: tint:657
Change-Id: I9efda26369c45c6f816bdaa53408d3909db403a1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53084
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Add test/intrinsics/intrinsics.wgsl.tmpl that generates a vast set of intrinsic overload permutations into test/intrinsics/gen/...
Add expected output for all of these, including 'SKIP' headers for those that currently fail.
Fixed: tint:832
Change-Id: Id6888df52c07f35e7a55199f2ad4d842c6e2595c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53051
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
This provides much more complete coverage than what we have in the
unit tests. We now test:
- All builtins, for all stages, both struct and non-struct
- Multiple location attributes for vertex and fragment stages, both
struct and non-struct
- Mixing builtins and location attributes, whilst mixing struct and
non-struct
- A few "interesting" cases of IO structs being shared between
different functions, stages, and with an SSBO variable
There are 7 skipped tests for MSL due to two different MSL bugs which
will be fixed in upcoming patches.
Change-Id: I8b802591762c8ff018e01bf37838551e353162b1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53120
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: James Price <jrprice@google.com>
Performs output validation with spirv-val for SPIR-V (as before), HLSL
validation with DXC, and MSL validation with the Metal Shader Compiler.
Disable HLSL tests that fail to validate
Bug: tint:812
Change-Id: If78c351b4e23c7fb50d333eacf9ee7cc81d18564
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51280
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This should be a u32, but we are missing validation rules for builtin
types (opened crbug.com/tint/861 for this). This meant that we were
generating invalid HLSL for this test.
Change-Id: Ib2189d98e5e515e41374372c0f8963df04eb1b01
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53001
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
This allows them to be used in various places that WGSL allows, such
as function return types and parameters, and as the type of the RHS of
an assignment.
Fixed: tint:814
Fixed: tint:820
Change-Id: Idb6a901b9a34e96bb9733cc158191e7b3bafaa0e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/52844
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>
Pull the HLSL transformation out to a standalone transform that can be
used by both HLSL and MSL.
The new E2E tests do not yet pass for MSL because they produce array
assignments, which will be addressed in the next patch.
Fixed: tint:826
Change-Id: Idc27c81ad45e3d4ab96d82663927d2fc1384618e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/52842
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>
Signed zeros are emitted.
Subormal numbers are emitted as hex float.
Handling Inf and NaN is unresolved in the spec.
See https://github.com/gpuweb/gpuweb/issues/1769
NaN tests are disabled, due to platform dependence.
Windows x86-64 seems to always set the high mantissa bit
on NaNs.
Fixed: tint:76
Change-Id: I06c69b93b04c869a63ec2f574022996acc7a6dbe
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/52321
Commit-Queue: David Neto <dneto@google.com>
Auto-Submit: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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>
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>
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>
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>
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>
Add a special-case for pointer-to-array types, where the * and the
variable name need to be enclosed in parentheses in between the array
element type and the size.
Move the `const` qualifier to before the array size.
Add E2E tests to cover all non-handle types used in various places.
Fixed: tint:822
Change-Id: I93b7d6867f92397aa47838ab2c94530b6e634617
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51823
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>
Much like arrays, the SPIR-V writer cannot cope with dynamic indexing of matrices.
Fixed: tint:825
Change-Id: Ia111f15e0cf6fbd441861a4b3455a33b82b692ab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51781
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
The special-case for zero-valued constructors is unnecessary, as an
empty initializer list already correctly zero-initializes for all
types. This was causing an additional {} to be emitted for empty
structures, which the MSL compiler rejects.
Fixed: tint:821
Change-Id: Ib48c73eadef15b517e14b248229ecfbbfeb13f81
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51822
Commit-Queue: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
BlockStatement for the root block of a function
Add some basic tests for this lot.
Bug: tint:812
Change-Id: I26b65717798cbff576a44bd78fbcb5c8f0d013e6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51368
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Alastair Donaldson <allydonaldson@googlemail.com>
Use this as part of the Spirv sanitizer.
Cleans up buggy dynamic array indexing logic in the SPIR-V writer.
Fixed: tint:824
Change-Id: Ia408e49bd808bc8dbf3a1897eb47f9b33b71fdfb
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51780
Commit-Queue: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
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>
A helper class and macro used to simplify scope-based assignment of
variables.
Change-Id: I02b3a05240a2c4628f813de931c40d8fba3cb07b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51480
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Change the type of the values in an ast::WorkgroupDecoration to be
ast::Expression nodes, so that they can represent both
ast::ScalarExpression (literal) and ast::IdentifierExpression
(module-scope constant).
The Resolver processes these nodes to produce a uint32_t for the
default value on each dimension, and captures a reference to the
module-scope constant if it is overridable (which will soon be used by
the inspector and backends).
The WGSL parser now uses `primary_expression` to parse arguments to
workgroup_size.
Also added some WorkgroupSize() helpers to ProgramBuilder.
Bug: tint:713
Change-Id: I44b7b0021b925c84f25f65e26dc7da6b19ede508
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51262
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This reverts commit 594075a2f0.
Reason for revert: This is not a complete solution. It does not cover assignment of variables of the same struct type, but of different access types. For example:
```
[[block]] struct S{ i : 32; };
var<storage> gr : [[access(read)]] S;
var<storage> gw : [[access(write)]] S;
fn f() {
var a : S;
a = gr;
gw = a;
}
```
With my CL, we currently generate invalid assignments because the new types of 'S' for 'gr' and 'gw' are now all different types. We would need to generate functions to assign between them.
In the end, we will drop this strategy, and instead, Dawn will be modified to not rely on names.
Original change's description:
> Add DuplicateStorageStruct transform
>
> This transform is currently required by Dawn for it's GLSL backend so
> that SPIRV-Cross does not end up renaming the structures internally,
> which it does to fix aliasing problems. We can remove this transform
> if/once Dawn using binding numbers rather than names for uniform/storage
> buffer data.
>
> Bug: tint:386
> Bug: tint:808
>
> CloneContext: allow insert before/after of a node marked for removal
> Change-Id: I3ff3b37bca62db07d5c759250dd4777e279b7a4f
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51403
> Commit-Queue: Ben Clayton <bclayton@google.com>
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Reviewed-by: Ben Clayton <bclayton@google.com>
TBR=bclayton@google.com,amaiorano@google.com,noreply+kokoro@google.com
Change-Id: I695347b0ffe8be23b9dc44885af378be56512406
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:386
Bug: tint:808
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51500
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
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>
James Price
David Neto
Ben Clayton
Corentin Wallez
Alastair Donaldson
Antonio Maiorano
Ryan Harrison
Brandon Jones
Sarah