Fixes#1642.
Since 1d textures cannot have mipmaps, MSL requires that the `level` argument to
texel accesses and dimension queries be a constexpr 0. For our purposes, just
omit the level argument altogether.
* glsl-in: Remove unneeded mutability from reference
* glsl-in: Fix composite constructors
In the recent rework of the constructors it seems that the logic for
composite types (arrays and structs) was accidentally removed by me.
* [hlsl-out] Write `mad` intrinsic for `fma` function
- This should be enough because we only support f32 for now.
- Adds a new test for WGSL functions, in the spirit of operators.wgsl.
- Closes#1579
* Add FMA feature to glsl backend
- I think this is right. Just iterate all known expressions in all
functions and entry points to locate any `fma` function call.
Should not need to walk the statement DAG.
* Transform GLSL fma function into an airthmetic expression when necessary
* Add tests for GLSL fma function tranformation
* Remove the hazard comment from the webgl test input
* Add helper method for fma function support checks
* Address review comment
* Add FindLsb / FindMsb
* Fixes and tests for FindLsb/FindMsb
* Add findLsb / findMsb as WGSL builtins
* Fix tests
* Fix incompatible type issue with MSL output
* Requested changes
* Test fewer cases of findLsb/findMsb
The IR doesn't allow having structs has bindings for entry point
input/output, so the glsl frontend must flatten them.
Glsl defines that the locations are sequential from the block base location.
The glsl spec defines that vector constructors flatten their arguments
and consume them in order, ignoring the remaining arguments if there are
more than needed
We can safely assume that deserialize is enabled for all output tests.
Supporting other modes of operation has low utility.
This change also adds an option to skip the explicit types for WGSL output.
Pointers should not be `DATA`: they can never be stored in anything. (Function
arguments are not storage; they're like `let` bindings.)
Un-`SIZED` values may only live in the `Storage` storage class, so creating
pointers to them in other storage classes is meaningless.
The `ARGUMENT` flag should be set only on pointers in those storage classes that
are permitted to be passed to functions.
See comments in code for details.
Fixes#1513.
The WGSL front end knows that globals in the `Handle` storage class do not
produce references, but it attempts to manage this in two places which trip over
each other. As a consequence, referring to a `let` binding holding a pointer to
a global flips the variable's type from WGSL `ptr` to WGSL reference, with
bewildering consequences.
* [spv-in] Change shadow.spv test input to use StorageBuffer.
The ecosystem around Naga will generally not be able to use Vulkan adapters that
don't support the SPV_KHR_storage_buffer_storage_class (which was incorporated
into SPIR-V 1.3), so we can assume it is present.
Changing the test not to use runtime-sized arrays in the Uniform storage class
will allow Naga to tighten up some validation checks.
* [spv-in] Permit pointers to runtime arrays only in StorageBuffer.
Fixes#1519.
The GLSL empty-global-name.frag test doesn't suffice because the GLSL front end
doesn't produce the same IR as the SPIR-V included in the bug report. As far as
I know, only a genuine SPIR-V input test can produce a global whose name is
`Some("")`.
Include the SPIR-V assembly source.
* Make default a switch case
Previously the default case of a switch statement was encoded as a block
in the statement but the wgsl spec defines it in such a way that the
default case ordering matters.
* [spv-out] Support for the new switch IR
* [dot-out] Use different labels for default cases
* Proof-of-concept for adding spans to validation errors.
Still missing: actually printing the damn stuff.
* Emit errors from analyzer in the CLI.
TODO: tests, I guess!
* Simplification refactoring: avoid avoiding allocations so vehemently.
* Mask helper traits with `as _`.
* Fix block iterator throwing up when span feature is disabled.
* Nest use statements.
* Add basic docs.
* Axe AddSpanResult.
Some compilers like shaderc introduce a full gl_PerVertex struct, this
includes gl_ClipDistance. Normally this isn't a problem since most
drivers optimize it away, but naga zero inits globals if they weren't
previously initialized. This causes gl_ClipDistance to be initialized to
zero which can be really bad for performance.
Minimize allocation in the namer. Heap-allocate string only when needed to
provide an owned key for a hash table, or to hold synthesized text. Try to reuse
allocations.
Many unnecessary trailing `_` characters are removed from test output. These
were all superfluous; separators are still inserted where necessary.
The original pointer access test used SPIR-V for its input because WGSL didn't
have a working pointer indirection operator at the time. Now that it does, we
can just write this test in WGSL directly.
Fixes#1432.
WGSL says:
> - The last member of the structure type defining the store type for a variable
> ... may be a runtime-sized array.
>
> - A runtime-sized array must not be used as the store type or contained within
> a store type in any other cases.
Thus, a struct whose final member is a struct whose final member is a
runtime-sized array is verboten.
Replace uses of `call_unique` with uses of `call` and `call_or`, which becomes
public. It's not clear when `call_unique` is correct to use, and avoiding a few
numeric suffixes here and there isn't worth it.
* Update WGSL grammar for pointer access.
Comes with a small test, which revealed a number of issues in the backends.
* Validate pointer arguments to functions to only have function/private/workgroup classes.
Comes with a small test. Also, "pointer-access.spv" test is temporarily disabled.
Automatically spills to a local variable function call arguments to
parameters expecting a pointer where the argument storage class isn't
function since the storage classes wouldn't match.
Treat expressions in `Function::named_expressions` like WGSL `let` declarations,
assuming that the Load Rule was applied to the rhs of the declaration, meaning
that their values are always `Indirection::Ordinary`.
Split `write_expr_plain_form` out from `write_expr_with_indirection`, to clean
up the parenthesis generation: no more `opened_paren` variable, just function
calls. This makes the early return for named expressions neater.
Fixes#1382.
Treat `TypeInner::ValuePointer` and `TypeInner::Pointer` as equivalent by
converting them to a canonical form before comparison.
Support `ValuePointer` in WGSL type output.
Fixes#1318.
Ensure that each distinct type occurs only once in `Module::types`, so that we
can use `Eq` on `Type` or `TypeInner` for type equivalence, without being
confused by differing `Handle<Type>` values that point to identical types.
This removes a number of duplicate types from the ir snapshots.
Fixes#1385.
Replace `Module::apply_common_default_interpolation` with a simpler function
that handles a single `Binding` at a time. In exchange for the simplicity, the
function must be called at each point function arguments, function results, and
struct members are prepared. (Any missed spots will be caught by the verifier.)
This approach no longer requires mutating types in the arena, a prerequisite for
properly handling type identity.
Applying defaults to struct members when the struct declaration is parsed does
have a disadvantage, compared to the old whole-module pass: at struct parse
time, we don't yet know which pipeline stages the struct will be used in. The
best we can do is apply defaults to anything with a `Location` binding. This
causes needless qualifiers to appear in some output. However, it seems that our
back end languages all tolerate such qualifiers.
Add support for float, vector and matrices targets.
Fix prefix and postfix being inverted (one was returning the value
of the other).
Remove an unneeded local indirection for prefix handling.
Add tests.
Fixes#1356.
Output for the test case in that issue is now:
error: the value indexed by a `[]` subscripting expression must not be a pointer
┌─ wgsl:5:14
│
5 │ let a = *pv[3]; // Problematic line
│ ^^ expression is a pointer
Could not parse WGSL
The old filenames were probably based on the shader stage, not the shader
function name, which the PR changed to do in response to a review suggestion.
`tests/out/glsl/operators.main.Fragment.glsl` is no longer generated because the
function in question changed to a Compute entry point.
`tests/out/wgsl/multiple_entry_points-glsl.wgsl` is no longer generated because
the input test was removed in 057dc310.
Make the parser code more closely follow the spec's grammar around
`unary_expression`, `postfix_expression`, and `singular_expression`.
Change the handling of postfix expressions (indexing, member/component access,
and swizzling) to apply the indirection at the appropriate time, resulting in
code improvements on all output formats. For example, where we used to generate
the following MSL:
metal::float4 _e13 = bar.matrix[3];
float b = _e13.x;
we now generate, simply:
float b = bar.matrix[3].x;
Propagate WGSL reference types correctly, so that parenthesizing expressions no
longer causes the Load Rule to be applied.
Together with #1332 (already landed), this is a replacement for #1312, and
unblocks #1352.
Fixes#1351.
Regression from ce676cf130.
We need to output (*d).mx rather than *d.mx, at least according to Tint.
wgsl-in seems to handle *d.mx just fine, which is likely a separate bug.
* [spv-in] New two pass parser based
* [spv-in] Allow expressions defined in dominant block in different scopes
* Make the patch non breaking
* [spv-in] Allow scope transfers in phi instructions
* [spv-in] Remove unused stuff
* [spv-in] Handle switch merges as breaks
* Remove no longer needed stuff
* Revert some changes to prepare to merge
* Remove dead code
* Don't spill into local if in scope
* [spv-in] Documentation, comments, some renaming for clarity.
* Address comments
Co-authored-by: Jim Blandy <jimb@red-bean.com>
* Fix GLSL output for non-fallthrough switch cases
Partially reverts 02c74b5002 and
fixes#1309.
* Fix indentation of control-flow WGSL code
* Clean up glsl-out switch case fallthrough handling
Only insert a break if needed
and if a case is fallthrough, insert a comment indicating this.
When we are leaning on robust buffer access to do the job for us, there's no
undefined behavior going on. So `UndefinedBehavior` suggests people are doing
something reckless even if they're not. The policy just says what Naga is doing,
and it shouldn't pretend to say what the rest of the system is doing.
* Implement lexing for all WGSL number literal types
* Move number literal test cases
* Adjust tests to match WGSL spec on number literals
Suffixes are not type names and currently only a plain `u` is supported
for uints. More specifically, `i` and `f` suffixes or suffixes with
widths in bits like `u32` are not supported at the moment.
* Add more tests for invalid number literal suffixes
* Replace code too new for Rust 1.43
* Implement parsing for hexadecimal integers
* Switch to enum number types, and Bytes for width
* Check for negative and leading zeros in int literals
* Implement parsing of hex floats with hexf-parse
* Update error message tests
* Update snapshot test output files
* Clean up lexer state machine code
* Clean up unexpected token error code
* Move number literal parsing to own submodule
Using 1D images should require either the `Sampled1D` or `Image1D` capability.
Using cube array images should require either the `SampledCubeArray` or
`ImageCubeArray` capability.
Introduce a helper type that manages the details of emitting SPIR-V
conditionals, tracking labels, branches, merge annotations and phi nodes.
Change index bounds checks to use this helper.
* [hlsl-out] flatten the entry point inputs
Previously, the logic was re-ordering the inputs according to the binding.
This breaks if one of the inputs is a struct. With this change, the struct
fields are also flattened into the fake entry point struct. We also
construct the original arguments at the beginning of the function.
* hlsl-out: completely separate the flattened IO structs from the original IR structs
Previously, we had heuristics to detect if a particular struct needs the fields
to be re-ordered. We'd re-order interface structs without layout, and the detection
was very fragile and easily wrong.
The new logic is spawning separate struct types if we need any re-ordering to happen.
It's solid, there are no heuristics.
Hlsl and wgsl don't support them directly so a polyfill is used taken
from the msl spec.
`asinh` = `log(x + sqrt(x * x + 1.0))`
`acosh` = `log(x + sqrt(x * x - 1.0))`
`atanh` = `0.5 * log((1.0 + x) / (1.0 – x))`
Some SPIR-V texture access instructions take coordinates as integers, others as
floats. The types of coordinates in Naga expressions generally match those in
SPIR-V, but Naga indices for arrayed textures are always integers, whereas
SPIR-V combines coordinates and array indices into a single vector, so indices
need to be cast to match the coordinate component type.
This commit makes `write_texture_coordinates` properly cast array indices to
match the coordinates' component type before combining them all into a single
result vector.
Fixes#1186.
In `back::spv::Writer::write_texture_coordinates`, OpCompositeConstruct can
concatenate scalars and vectors, so when combining coordinates with an array
index, there is no need to extract the coordinate vector's components
individually: once the index has been converted to the appropriate component
type, it can be combined with the coordinates in a single instruction.
The new api allows for reuse while keeping some allocations and to
please the borrow checker in future work, it also splits the parser into
logical modules to make it easier to work on.
Glsl defines that if a matrix constructor is called with only a scalar
it's result will be the matrix where the scalar is used if the row and
column are equal and everywhere else a zero
Glsl allows things such as
```glsl
layout (location = 0) in vec4 a;
vec4 b = a;
```
In this case `b` depends on `a` which can't be known at compile time,
so it's necessary to add a small prologue to the entry point to
initialize it.
* Resurrect texture_storage_* tests
* Test parsing of `var<storage,write>`
* Default storage textures to READ
* Restore default features
* Fix glsl/hlsl/msl/spv front and back ends
* Add missing test outputs
* All-around fixes for the storage access
Co-authored-by: Dzmitry Malyshau <kvarkus@gmail.com>
* Hack in support for PrimitiveID on Vulkan
* Rename to PrimitiveIndex and add preliminary support for GLSL, HLSL and MSL
* Implement as an extra WGSL built-in
* Update extra.wgsl outputs
* Run rustfmt; fix WGSL writer
* Add rustfmt changes I forgot
* Update extra.wgsl test output for WGSL fix
* Bump macOS version to 10.15 in validate-msl to support primitive_index
WGSL will require this. Note that this still might
cause some issues with multi-entry-point GLSL
that I didn't know how to handle. That is, we will
handle unused builtin inputs but not unused
builtin outputs correctly right now. This is an existing
issue though, not a regression.
This is also provided as an option, but I feel like
the more correct approach is to never strip linkage
variables. We'll see though.
