Migrate some of the validation logic over to use the results of behavior
analysis.
The most significant changes are:
* Unreachable-statements now consider merge-points of control flow. For
example, if all branches of a if-statement or switch-statement either
return or discard, the next statement will be considered unreachable.
* Unreachable statements are no longer an error, but a warning. See
https://github.com/gpuweb/gpuweb/issues/2378.
* Statements that follow a loops that does not break, or have a
conditional will now be considered unreachable.
* Unreachable statements produced by the SPIR-V reader are now removed
using the new RemoveUnreachableStatements transform.
Some other new changes include additional validation for the continuing
block for for-loops, to match the rules of a loop continuing block.
The new cases this validation is testing for are not expressible in
WGSL, but some transforms may produce complex continuing statements that
might violate these rules. All the writers are able to decay these
complex for-loop continuing statements to regular loops.
Bug: tint:1302
Change-Id: I0d8a48c73d5d5c30a1cddf92cc3383a692a58e61
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/71500
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
This change implements the behavior analysis for expressions and
statements as described in:
https://www.w3.org/TR/WGSL/#behaviors
This CL makes no changes to the validation rules. This will be done as a
followup change.
Bug: tint:1302
Change-Id: If0a251a7982ea15ff5d93b54a5cc5ed03ba60608
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68408
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
This new expectation file was added with a similar parent that changed the expected output.
Bug: tint:1112
Change-Id: Icb6175f68b6459385933d39e07460eb863ecb9ff
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/71101
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Now that I've landed this change to Dawn to disable FXC optimizations:
https://dawn-review.googlesource.com/c/dawn/+/70700,, we can reland this
change. The Tint-into-Dawn roll was failing because FXC would miscompile
certain loops into infinite loops when not unrolled.
Also reland
test/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.hlsl
Bug: tint:1112
Bug: dawn:1203
Change-Id: I641d68864b833e0fbe3b117d397b89ae96482536
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/71000
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
If we encounter a global twice, the stack vector was not popped, which could trigger false cyclic-dependency errors. Because the cyclic dependency did not actually exist, later logic could break, expecting to find a global->global edge which did not exist.
Fixed: chromium:1273451
Change-Id: I9d99f1aeeaea042d9ed847a878c4717803122240
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70980
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
It's surprising that the fuzzers didn't find this, but I've confirmed that at least one of the fuzzers will fail with this case.
By adding a e2e test case, this will be used as a seed for the fuzzer corpus.
Bug: tint:1321
Change-Id: Ibf4bfa50bf376ae4cba401dfc31a2498fa3b9eec
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70982
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
RemovePhonies was transforming:
_ = f32(1)
_ = vec2<f32>(1.0, 2.0)
into:
f32(1)
vec2<f32>(1.0, 2.0)
Which the resolver gets grumpy about, as these are expressions, not statements.
Fixed: chromium:1273230
Change-Id: Ie85d3cee705fa3f792db686c021d76331e241f17
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70960
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
BUG=tint:1019
Change-Id: Ia462080877a97348c5589bfa71231a832a7ebfd3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70081
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Add transform::Unshadow to renamed shadowed symbols. Required by a
number of other transforms.
Replace Resolver symbol resolution with dep-graph.
The dependency graph now performs full symbol resolution before the
regular resolver pass.
Make use of this instead of duplicating the effort.
Simplfies code, and actually performs variable shadowing consistently.
Fixed: tint:819
Bug: tint:1266
Change-Id: I595d1812aebe1d79d2d32e724ff90de36e74cf4b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70523
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
They are always used together, and can actually be simplified when
combined.
This transform cannot currently deal with shadowing, and will need
need to depend on a new 'unshadow' transform.
Fixes a long-standing bug where we'd get an ICE if we attempt to inline
a pointer let declaration with side-effects in a for-loop initializer.
Fixed: tint:1321
Bug: tint:819
Change-Id: I236fed688e33a4996e47310b5ece44c991b5249f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70661
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
These have all been removed from the spec as we cannot guarantee they will behave as expected on the various backends.
Bug: tint:1312
Change-Id: I9d7d81cfdc44489fffe08c5183ed8da84901a024
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70665
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Looks like this should have been part of bc5d8f6b
Bug: tint:1112
Change-Id: I62e98bbb5f8ca29a171d5d5baa65af55b3a090d4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70664
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
A find or return-default utility.
Rename `find_or_replace.h` to `map.h` which contains both these utilties that operate on maps.
Change-Id: Iaa76ea3f5c5a1210e413c131dd0556c126ee5d0a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70521
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
The storageBarrier() builtin causes flushes to be visible to
all invocations within a workgroup, not the whole device.
The limitation is inherited from Metal.
Fixes: tint:1310
Change-Id: I6f94faa88bd3c7b6cec0601312c6c65a907c5973
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69800
Auto-Submit: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This reverts commit 11d09f2fe7.
Reason for revert: Failing roll of Tint to Dawn: https://dawn-review.googlesource.com/c/dawn/+/70100
Original change's description:
> HLSL: force FXC to never unroll loops
>
> Emit the "[loop]" attribute on "for" and "while" so that FXC does not
> attempt to unroll them. This is to work around an FXC bug where it fails
> to unroll loops with gradient operations.
>
> FXC ostensibly unrolls such loops because gradient operations require
> uniform control flow, and loops that have varying iterations may
> possibly not be uniform. Tint will eventually validate that control flow
> is indeed uniform, so forcing FXC to avoid unrolling in these cases
> should be fine.
>
> Bug: tint:1112
> Change-Id: I10077f8b62fbbb230a0003f3864c75a8fe0e1d18
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69880
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Reviewed-by: Ben Clayton <bclayton@google.com>
> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
# Not skipping CQ checks because original CL landed > 1 day ago.
Bug: tint:1112
Change-Id: I8e8f3c0abfa6e6bc5d0e67af9428a46ef867d5c1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70540
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>
No need for the HLSL-style repeated swizzle; GLSL allows construction
of a vector from a scalar value of the component type.
Bug: tint:1317
Change-Id: Ia0afe3012cbb56716a2d1c5c3849dd662a5ff89c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70342
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
'__' is reserved in C++, and the 'match__' and 'build__' functions are causing OSS-fuzz builds to fail.
Add the change in tint behavior to the OT notes.
Add end to end tests for underscores. While the GLSL and MSL compilers seem to accept leading and double underscores in identifiers, the tint build failure has highlighted we have more work to do here (crbug.com/tint/1319)
Fixed: oss-fuzz:41214
Bug: tint:1292
Bug: tint:1319
Change-Id: I32b7bf4e0cff26e678b788457f90452c2503da50
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70480
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Performs a module-scope (global) declaration dependency analysis, so
that out-of-order global declarations can be re-ordered into dependency
order for consumption by the resolver.
The WGSL working group are currently debating whether out-of-order
declarations should be included in WebGPU V1, so this implementation
currently errors if module-scope declarations are declared out-of-order,
and the resolver does not currently use this sorted global list.
The analysis does however provide significantly better error diagnostics
when cyclic dependencies are formed, and when globals are declared
out-of-order.
The DependencyGraph also correctly now detects symbol collisions between
functions and types (tint:1308).
With this change, validation is duplicated between the DependencyGraph
and the Resolver. The now-unreachable validation will be removed from
the Resolver with a followup change.
Fixed: tint:1308
Bug: tint:1266
Change-Id: I809c23a069a86cf429f5ec8ef3ad9a98246766ab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69381
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Try to lex identifiers before punctuation, and backtrack if we only
see a single underscore or something that starts with two underscores.
Rename the modf and frexp return types to start with two underscores.
Spec PR:
https://github.com/gpuweb/gpuweb/pull/2326
Change-Id: Id283af100babfe84faa183345cb8a60848140caa
Fixed: tint:1292
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70160
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Use imageSize() on images, not textureSize().
In GLSL, the LOD parameter to textureSize() is mandatory for
sampled textures, so emit a default 0 if not supplied. (Also, don't pack
the level into the coords argument; that's an HLSLism.)
GLSL returns the array size of array textures in the final component
of textureSize(); remove it for WGSL.
Write the subtype of storage images correctly (uimage*, iimage*, etc).
This required a bit of cleanup to move "writeonly" ahead of subtype
emission.
Bug: tint:1298
Change-Id: Ica1cec0f833a9b684143c8b0cf6d090fb511a7d2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70140
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
ArrayLengthFromUniform is needed for correct bounds checks on
dynamic storage buffers on D3D12. The intrinsic GetDimensions does
not return the actual size of the buffer binding.
ArrayLengthFromUniform is updated to output the indices of the
uniform buffer that are statically used. This allows Dawn to minimize
the amount of data needed to upload into the uniform buffer.
These output indices are returned on the HLSL/MSL generator result.
ArrayLengthFromUniform is also updated to allow only some of the
arrayLength calls to be replaced with uniform buffer loads. For HLSL
output, the remaining arrayLength computations will continue to use
GetDimensions(). For MSL, it is invalid to not specify an index into
the uniform buffer for all storage buffers.
After Dawn is updated to use the array_length_from_uniform option in the
Metal backend, the buffer_size_ubo_index member for MSL output may be
removed.
Bug: dawn:429
Change-Id: I9da4ec4a20882e9f1bfa5bb026725d72529eff26
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69301
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Austin Eng <enga@chromium.org>
Emit the "[loop]" attribute on "for" and "while" so that FXC does not
attempt to unroll them. This is to work around an FXC bug where it fails
to unroll loops with gradient operations.
FXC ostensibly unrolls such loops because gradient operations require
uniform control flow, and loops that have varying iterations may
possibly not be uniform. Tint will eventually validate that control flow
is indeed uniform, so forcing FXC to avoid unrolling in these cases
should be fine.
Bug: tint:1112
Change-Id: I10077f8b62fbbb230a0003f3864c75a8fe0e1d18
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69880
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Use uintBitsToFloat.
Bug: tint:1306
Change-Id: Ie9a5e14c13c0d63b57c126f16c4e2a5c7a77e3f7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69740
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Looks like a typo in the test cases I wrote.
Change-Id: Ieb4d8ce28827e47ab0baef7b1178395d97f90ace
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69841
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Implements MultiplanarExternalTextureTransform to allow transforming a
texture_external binding into two texture_2d<f32> bindings and a uniform
buffer binding. Transforms textureSampleLevel and textureLoad calls with
a texture_external parameter into custom functions that can handle both
single-plane RGBA or bi-planar YUV. Includes tests.
Bug: dawn:1082
Change-Id: Icb6d8b0f3773feca01c833171f07230c3531f3aa
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68620
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
In GLSL, all identifiers beginning with gl_ are reserved (not just those
explicitly named in the spec), and so any found in WGSL must be renamed.
Bug: tint:1304
Change-Id: I92ed7ec674620f67775378ecb8debcfdb4b5bbb4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69701
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
In GLSL, relational operators are only valid for scalar operands. For
vector operands, lessThan, greaterThan, etc must be used.
Bug: tint:1303
Change-Id: Ia800f89111630c756dc1b30ef0c6858fb520fb16
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69561
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
In an attempt to better follow the guidelines in docs/end-to-end-tests.md.
Change-Id: Id012004e8a376598bb69cf364ca7ca0735632a2e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69342
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Add a new 'Target' to the ast::CallExpression, which can be either an
Identifier or Type. The Identifier may resolve to a Type, if the Type is
a structure or alias.
The Resolver now resolves the CallExpression target to one of the
following sem::CallTargets:
* sem::Function
* sem::Intrinsic
* sem::TypeConstructor
* sem::TypeCast
This change will allow us to remove the type tracking logic from the WGSL
parser, which is required for out-of-order module scope declarations.
Bug: tint:888
Bug: tint:1266
Change-Id: I696f117115a50981fd5c102a0d7764641bb755dd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68525
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
An element-wise vector transformation utility function.
Similar to JS/TS's map() method on arrays.
Change-Id: I4baf52daa918f2e7bf5f9b4af13894fe66826f7c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69103
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
The 'tint_unittests_src' source_set needed to depend on gmock_and_gtest,
which is normally done with the 'tint_unittests_source_set' template,
but this is the root source_set for the unit tests.
To fix, these root-level sources have been moved to
'tint_unittests_core_src' which uses the 'tint_unittests_source_set'
template.
Change-Id: I0d76c55c744646fd7deff22406a668f8c6cdf519
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69220
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Sarah Mashayekhi <sarahmashay@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Prevents errors when we have the same _test.cc file in two different
directories.
Change-Id: I62eaea9452762670b7a24cdb2d7b0bef4fe52280
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/69102
Kokoro: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
To fix this, we trick the compiler by wrapping the function body with an
if (true) { <function body> } followed by returning an unused value of
the return type.
Bug: tint:1081
Change-Id: I763bf768f40d07a1045f0a70017bb40d488c8428
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68822
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
The object is not always an array. The index can be applied to vectors
too.
Change-Id: Ifb63d1862090d28cb48d692870e9dd01ddbce5df
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68841
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Literals are now expressions, so in keeping with all the other
expression types, suffix the class name with Expression.
I'm not overly keen on requiring everything to have an Expression
suffix, but consistency is better than personal preference.
Bug: tint:888
Change-Id: Ida1f1b98c71d5d0fe6399bca938010cd5d3a00c3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68840
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Just make Literal an expression. The ScalarConstructorExpression
provides no real value, aside from having a ConstructorExpression base
class that's common between ScalarConstructorExpression and
TypeConstructorExpression. TypeConstructorExpression will be folded into
CallExpression, so this hierarchy will serve no purpose.
First step in resolving the parser ambiguity of type-constructors vs
type-casts vs function calls.
Bug: tint:888
Change-Id: I2585d5ddbf6c0619a8f24c503e61ebf27c182ebe
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68524
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Some pass DXC, but not FXC (e.g. continue_in_switch).
Some now pass (e.g intrinsics/gen/atan2/*.hlsl)
Some now fail gracefully instead of asserting (e.g. 807.spv)
Change-Id: I92b17fcadc7850de5bd93ff07507cea7e5487fc9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68820
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
In HLSL code, if a switch statement has only a default case, FXC will
effectively ignore the code in that case. In this change, we detect this
and work around it by emitting the code in the default block without the
switch.
Bug: tint:1188
Change-Id: I69b405cdb4c669fb093eb49aa138923419dcf8f8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68440
Kokoro: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
The semantic nodes cannot be fully immutable, as they contain cyclic
references. Remove Resolver::CreateSemanticNodes(), and instead
construct and mutate the semantic nodes in the single traversal pass.
Give up on trying to maintain the 'authored' type names (aliased names).
These are a nightmare to maintain, and provided limited use.
Significantly simplfies the Resolver, and allows us to generate more
semantic to semantic references, reducing sem -> ast -> sem hops.
Note: This change introduces constant value propagation across constant
variables. This is unlocked by the earlier construction of the
sem::Variable.
Change-Id: I592092fdc47fe24d30e512952511c9ab7c16d7a1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68406
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Update or remove tests that try to do this.
Fixed: tint:491
Change-Id: I1f351a4abf68ae9bc6b100885fb1bcea08b31211
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68242
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
WGSL does not support pointer-to-vector-component, so the SPIR-V
reader needs to sink these pointers into their use.
Bug: tint:491
Change-Id: Ib5ae87d2f6bbac13280314ba11369d7ced505b56
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68241
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
This removes a lot of awkward logic from the MSL writer, and means
that we now handle all module-scope variables with the same transform.
Change-Id: I782e36a4b88dafbc3f8364f7caa7f95c6ae3f5f1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67643
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
These expressions were not emitted as spec-ops, failing new CTS tests.
Change-Id: I56e8f56a22de2b8dac9a8bd7a2d694d8d81dca35
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67480
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Use a transform to convert these to the vector form for the MSL and
SPIR-V backends.
MSL only has the scalar form from version 2.0 onwards.
Fixed: tint:1123
Change-Id: I384abd9872d9eae52a10a37cbd6aa96004692e9c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67360
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Use spec-spelling of 'constructible'.
Add missing test file to 'test/BUILD.gn'
See https://dawn-review.googlesource.com/c/tint/+/67064
Change-Id: Ie39773617fd0a363d63cc6449bf3905c9eb6786d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67380
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>
It's been removed from the spec:
https://github.com/gpuweb/gpuweb/pull/2127
Fixed: tint:1213
Change-Id: I163fe807765bb1ac0580b398f4897daea555216a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67067
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
The `Ignore()` intrinsic is about to be deprecated, so don't use it for testing.
Bug: tint:1213
Change-Id: I314ecaeb9a9c337c7b6980189054120a74807ebd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67066
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Produce a warning if the attribute is missing for integral vertex
outputs or fragment inputs. This will become an error in the future,
as per the WGSL spec.
Add the attribute to E2E tests.
Bug: tint:1224
Change-Id: Ia1f353f36cb7db516cf9e8b4877423dec3b3e711
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67160
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>
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>
The intrinsics that did anything useful with this were deprecated
several releases ago.
Change-Id: I79e3c901b6a78583853a067ec46cfa98e346517c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/66262
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
These have been deprecated for multiple chrome releases.
Change-Id: I4cc05a74ff8f085e6d13f93aefb93077480e52f5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/66261
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This is a modified version of the HLSL writer.
Basic types, arrays, entry points, reserved keywords, uniforms,
builtin uniforms, structs, some builtin functions, zero initialization
are implemented. Textures, SSBOs and storage textures in particular are
unimplemented. All the unit tests "pass", but the output is not correct
in many cases.
triangle.wgsl outputs correct vertex and fragment shaders that pass
GLSL validation via glslang. compute_boids.wgsl outputs a valid but not
correct compute shader.
Change-Id: I96c7aaf60cf2d4237e45d732e5f51b345aea0552
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57780
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
GLSLstd450SAbs expects a *signed* integer.
abs() of an unsigned number is now a no-op.
Fixes WebGPU CTS tests:
webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:*
Bug: tint:1194
Change-Id: I65c5e9f2f03aac0b788b9ba88c383cbec136d7c6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65620
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
MSL has a limit on the number of threadgroup memory arguments, so use
a struct to support an arbitrary number of workgroup variables.
This commit introduces a `State` object to this transform, which is
used to track which structs have been cloned eagerly, in order to
avoid duplicating them.
Bug: tint:938
Change-Id: Ia467db186e176a08f160455eab5fd3b3662f56b8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65360
Auto-Submit: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Correctly rename fields when combining two or more input structures together into a single input structure.
Bug: chromium:1251009
Change-Id: I0c7ab5ed3116b97035e100d1ef96e772e819f640
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/64545
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Catch up to include all tests emitted by tint_unittests --dump-spirv
Change-Id: Ia4d5f75782bcc62f8cfeb0615a942e95e563289d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65041
Auto-Submit: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
This reverts commit af8cd3b7f5.
Reason for revert: breaking roll into Dawn.
Original change's description:
> msl: Use a struct for threadgroup memory arguments
>
> MSL has a limit on the number of threadgroup memory arguments, so use
> a struct to support an arbitrary number of workgroup variables.
>
> Bug: tint:938
> Change-Id: I40e4a8d99bc4ae074010479a56e13e2e0acdded3
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/64380
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Auto-Submit: James Price <jrprice@google.com>
> Reviewed-by: Ben Clayton <bclayton@google.com>
> Commit-Queue: James Price <jrprice@google.com>
TBR=bclayton@google.com,jrprice@google.com,noreply+kokoro@google.com,tint-scoped@luci-project-accounts.iam.gserviceaccount.com
Change-Id: I58a07c4ab7e92bda205e2bbbab41e0b347aeb1e8
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:938
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65162
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
Kokoro: Corentin Wallez <cwallez@chromium.org>
MSL has a limit on the number of threadgroup memory arguments, so use
a struct to support an arbitrary number of workgroup variables.
Bug: tint:938
Change-Id: I40e4a8d99bc4ae074010479a56e13e2e0acdded3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/64380
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Update one test to avoid this error.
Fixed: tint:1179
Change-Id: Id41b0eb0f404648de4e86a835fe43f1729cb4696
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/64464
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>
Use a threadgroup memory argument for any workgroup variable that
contains a matrix.
The generator now provides a list of threadgroup memory arguments for
each entry point, so that the runtime knows how many bytes to allocate
for each argument.
Bug: tint:938
Change-Id: Ia4af33cd6a44c4f74258793443eb737c2931f5eb
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/64042
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
For HLSL, use the new NumWorkgroupsFromUniform transform, and expose
the binding point to use for the generated uniform as a backend
option.
The MSL mapping is trivial, and it was already implemented for WGSL
and SPIR-V.
Bug: tint:752
Change-Id: I4bd37b5d26181629d72b152fe064a60caf8ecdc5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63962
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This transform scans entry points for struct parameters that contain
the num_workgroups builtin, and replace accesses to these members with
a value loaded from a uniform buffer.
This will be used by the HLSL backend to implement the num_workgroups
builtin.
Bug: tint:752
Change-Id: Iefab3b14af8a08a6135348fded368a06d932e915
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63961
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Invoke the required transforms directly in the SPIR-V backend.
Change-Id: I78dc667d5c4c9c1d4da13ef5a99ece831c103982
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63801
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Invoke the required transforms directly in the HLSL backend.
Change-Id: I9465fef375dd4dad6a91c1e7e16ede6401b9bfc0
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63800
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Invoke the required transforms directly in the MSL backend.
Change-Id: Id8026b1a64415fbe363f8f8a5790e8216cd12c68
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63620
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
This CL changes the way that the resolver traverses expressions to avoid stack overflows for deeply nested expressions.
Instead of having the expression resolver methods call back into
Expression(), add a TraverseExpressions() method that collects all the
expression nodes with a simple DFS.
This currently only changes the way that Expressions are traversed. We
may need to do the same for statements.
Bug: chromium:1246375
Change-Id: Ie81905da1b790b6dd1df9f1ac42e06593d397c21
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63700
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This is the HandleModuleScopeVars() part of the MSL sanitizer moved
verbatim to a standalone transform. The transform code is unchanged,
but some expected test outputs are different as this is now tested in
isolation instead of along with the rest of the sanitizer transforms.
This is step towards removing the sanitizers completely.
Change-Id: I7be826e2119451fc2ce2891740cc94f978e7d5a1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63583
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Use this from the HLSL and SPIR-V sanitizers, instead of duplicating
this logic for them.
This is step towards removing the sanitizers completely.
Change-Id: Ifa9f23d84fd3505d30a928c260181a699c5f1783
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63582
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
No new tests or any changes, just moving to a separate file in
preparation for adding lots more of these tests soon.
Change-Id: Iaa7eef52384e702c395a6db312fef19e22507644
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63580
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: 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>
Polyfill this for HLSL using an atomic add with the operand negated.
Fixed: tint:1130
Change-Id: Ifa32d58973f1b48593ec0f6320f47f4358a5a3a9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62760
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>
For these tests, we only really care that we can successfully consume
them and generate valid output for each backend. Having the expected
files in the tree generates significant churn for any change to how we
generate backend code, which makes it hard to inspect diffs.
Change-Id: Ic98c248081144c0fb1791f1303eaf6d459548e3d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62720
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
These operators are not defined in the metal namespace when the vector
operands are packed.
Fixed: tint:1121
Change-Id: I2e8f4302e08117ca41bac6c05fb24a70d1215740
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62480
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
MSL vectors with other widths already match WGSL's rules for alignment
and size.
Change-Id: I237052372463ea8323eab47c3b4ca90c6d8afcc3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62600
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Remap all resources into a flat namespace, to allow tests to pass when
multiple resources use the same binding number.
Fixed: tint:959
Change-Id: I58ed07c789e1ea90fc370ceba73b9d8292902549
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/61261
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
The location decoration in the test source was recently changed to fit
in MSL's allowed range.
Change-Id: Ia8ff2cc453f5af3ed084ede42546ef6750bd27a9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/61260
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This new test was added around the same time as the entry point IO
rework, so the (now incorrect) test output made it past the bots.
Change-Id: I89fc4041b9cd00cd363ba61d07371554263eca96
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/61460
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@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>
FXC errors on these, and they are undefined behavior in WGSL.
Bug: tint:1083
Change-Id: I7643fdc6991f8729f274535b603b761398412398
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60500
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
FXC does not support fallthrough case statements (DXC does).
Fixed: tint:1082
Change-Id: I82e1add5455e438056259f773f34bf9db05970b4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60480
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Add the unit tests samples from src/reader/spirv when:
- they are valid for Vulkan 1.0 (plus some common extensions)
- they should translate to valid WGSL
Bug: tint:1043
Change-Id: I40a01990dbc40aff5cf7ace0b1aabfd0e437f638
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60000
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: David Neto <dneto@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
https://dawn-review.googlesource.com/c/tint/+/60213 special cased ignore() to work around tint:1046.
This fix produced bad output for structures when they are fully decomposed into ByteAddressBuffers, as the final HLSL references a structure that no longer exists.
Fixes CTS tests, and tint->dawn roll.
Change-Id: If6eab083c5f0bcca4a90c582df255b77e97a8e9f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60347
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Spread the array zeroing across as many workgroup invocations as possible.
Bug: tint:910
Change-Id: I1cb5a6aaafd2a0a4093ea3b9797c173378bc5605
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60203
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
This reverts commit e5dbe24e94.
Reason for revert: Makes the Tint-Dawn roll fails because of
MSL compilation errors on as_type<uint>(-2147483648):
as_type cast from 'long' to 'uint' (aka 'unsigned int') is not allowed
as_type<uint>(-2147483647) compiles fine, so this is most
likely because the MSL compiler types the literal as a long
(since without the - it is larger than the max int32).
Original change's description:
> 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>
TBR=dneto@google.com,bclayton@google.com,jrprice@google.com,amaiorano@google.com,noreply+kokoro@google.com,tint-scoped@luci-project-accounts.iam.gserviceaccount.com
Change-Id: I3e3384a9185013bb141a1b7b9b22bad8571bbc50
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:124
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60345
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Auto-Submit: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
For loops only support assignments or function calls for the continuing statement.
Fixed: tint:1064
Change-Id: I07065b2119e7b9f97ca7e46b1464fd72333ca429
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60212
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@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>
The second parameter must not be a u32.
Fixed: tint:1078
Bug: tint:1079
Change-Id: Id7a9cd881c4fec0f262931c2e4c263310e59c25d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60204
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@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>
This code was implicitly assuming that all resources it was looking
for would be directly referenced at the intrinsic callsite, and not
passed via function parameters.
This was causing a crash in more complex cases.
The inspector code has been updated to handle cases where the
resources are not being directly referenced.
Unneeded calls to GenerateSamplerTargets() are removed.
Utility function GetOriginatingResources() is added to handle walking up
call sites to resolve resources.
Text shader based test runner is added to the Inspector tests to make
expressing complex tests easier.
BUG=tint:967
Change-Id: I2ecb6d57c518003da59f38b261bae4d62ce7e6ac
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59340
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
The continuing block can exit the loop in very constrained ways:
When a break statement is placed such that it would exit from a loop’s
§ 7.3.8 Continuing Statement, then:
- The break statement must appear as either:
- The only statement in the if clause of an if statement that has:
- no else clause or an empty else clause
- no elseif clauses
- The only statement in the else clause of an if statement that has an
empty if clause and no elseif clauses.
- That if statement must appear last in the continuing clause.
By design, this allows a lossless round-trip from SPIR-V to WGSL and
back to SPIR-V. But that requires this special case construct in WGSL
to be translated to an OpBranchConditional with one target being
the loop's megre block (which is where 'break' branches to), and the
other targets the loop header (which is the loop backedge). That
OpBranchConditional takes the place of the normal case of an
unconditional backedge.
Avoids errors like this:
continue construct with the continue target X is not
post dominated by the back-edge block Y
Fixed: 1034
Change-Id: If472a179380b8d77af746a3cd8e279c8a5e56b37
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59800
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>
Add `transform::DecomposeStridedMatrix`, which replaces matrix members of storage or uniform buffer structures, that have a [[stride]] decoration, into an array
of N column vectors.
This is required to correctly handle `mat2x2` matrices in UBOs, as std140 rules will expect a default stride of 16 bytes, when in WGSL the default structure layout expects a stride of 8 bytes.
Bug: tint:1047
Change-Id: If5ca3c6ec087bbc1ac31a8d9a657b99bf34042a4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59840
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Implemented for all readers and writers.
Cleaned up some verbose code in sem::Function and the Inspector in the
process.
Fixed: tint:1032
Change-Id: Ia6f2f59e6d2e511c89160b97be990e8b7c9828d9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59664
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>
This forbids let declarations from having handle types.
Change-Id: I6f7467b0fa3963711ec705e1a81bfdd2c550feee
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59801
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Several more tests are passing now that we have hex float parsing and
vector bitcasts in the HLSL backend.
Change-Id: I2809c83aa78afa7cfec187a2cb1671f79e06a876
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59822
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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>
Ben Clayton
Antonio Maiorano
Ryan Harrison
David Neto
Stephen White
James Price
Austin Eng
Brandon Jones
Stephen White
Corentin Wallez
Alastair Donaldson
Sarah