Metal 1.x does not support swizzling on packed_vec types.
Use array-index for single element selection (permitted on LHS and RHS of assignment)
Cast the packed_vec to a vec for multiple element swizzles (not permitted as the LHS of an assignment).
Fixed: tint:1249
Change-Id: I70cbb0c22a935b06b3905d24484bdc2edfb95fc2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67060
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Change ast::Array to use an ast::Expression for its `size` field. The
WGSL frontend now parses the array size as an `primary_expression`,
and the Resolver is responsible for validating the expression is a
signed or unsigned integer, and either a literal or a non-overridable
module-scope constant.
The Resolver evaluates the constant value of the size expression, and
so the resolved sem::Array type still has a constant size as before.
Fixed: tint:1068
Fixed: tint:1117
Change-Id: Icfa141482ea1e47ea8c21a25e9eb48221f176e9a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63061
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>
In https://dawn-review.googlesource.com/c/tint/+/62444 the Resolver validated that there are no parameters of the same function with the same name, but this also introduced validation that errors if parameters shadow a module-scope variable.
The WGSL spec allows for shadowing, but Tint so far has not implemented this support.
There are transforms that generate functions that presume parameter <-> module-scope variable shadowing is okay. DecomposeMemoryAccess is one of these.
This fixes those transforms which could generate programs that fail validation.
Bug: chromium:1242330
Fixed: tint:1136
Change-Id: Id6ec59bbdb398b3b2a23312115a7c1dadf433e98
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62900
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
SPIR-V spec states:
> Each structure-type member that is a matrix or array-of-matrices must have be decorated with a MatrixStride Decoration
As already pointed out in https://dawn-review.googlesource.com/c/tint/+/59840, we were not handling arrays-of-matrices.
To do this correctly, we need the ast::StrideDecoration to be placed on the Matrix type, which is a much bigger change to support.
For now, chase the type, and error if we have a custom MatrixStride on an array of matrices, otherwise drop the decoration.
Bug: tint:1049
Fixed: tint:1088
Change-Id: Idcb75b3df88040836a03a14e0ca402ebee7be9a7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60923
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>
The refactored CanonicalizeEntryPointIO transform makes it much easier
to handle SPIR-V style IO as well, and doing this removes a lot of
duplicated code. Remove all of the SPIR-V transform code for shader IO
and vertex point size.
Bug: tint:920
Change-Id: Id1b97517619b4d2fd09b45d5aee848259f3dfa77
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60840
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 is a major reworking of this transform. The old transform code
was getting unwieldy, with part of the complication coming from the
handling of multiple return statements. By generating a wrapper
function instead, we can avoid a lot of this complexity.
The original entry point function is stripped of all shader IO
attributes (as well as `stage` and `workgroup_size`), but the body is
left unmodified. A new entry point wrapper function is introduced
which calls the original function, packing/unpacking the shader inputs
as necessary, and propagates the result to the corresponding shader
outputs.
The new code has been refactored to use a state object with the
different parts of the transform split into separate functions, which
makes it much more manageable.
Fixed: tint:1076
Bug: tint:920
Change-Id: I3490a0ea7a3509a4e198ce730e476516649d8d96
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60521
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>
Fix the Renamer to preserve builtin structure member names.
Fix the HLSL writer to emit the modf / frexp result type even if there is no private / function storage usage of the types.
Fixed: chromium:1236161
Change-Id: I93b9d92980682f9a9cb090d07b04e4c3f6a2f705
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60922
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
The new vk-gl-cts tests have uncovered a whole bunch of FXC issues,
which have been filed as tint bugs.
Bug: tint:998
Bug: tint:1080
Bug: tint:1038
Bug: tint:1081
Bug: tint:1082
Bug: tint:1083
Change-Id: I0d14370f94647dfd9c7088e0b782c3b415c78ee7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60211
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
When building a vector via tint::writer::AppendVector, and the
vector argument is already a vector constructor, expand that
vector constructor into its components only when those components
are all scalars. This avoids a type breakage which can occur with cases
like this:
vector argument is:
vec2<i32>(vec2<u32>(0u,1u))
scalar argument is:
2
Before this fix, the result was:
vec2<i32>(0u, 1u, 2);
But should be this instead:
vec3<i32>(vec2<u32>(0u,1u),2)
This was noticed in SPIR-V writer output when forming a coordinate
vector from a an unsigned WGSL coordinate vector with a signed array
vector.
Fixed: tint:1048
Change-Id: Id46665739cc23da0ca58b9baabf7b4531b86350b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60040
Auto-Submit: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: 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>
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 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>
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>
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>
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>
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>
