For simplicity's sake our initial implementation of external textures
will not support binding arrays of external textures. We should
therefore reject any shaders which use them during validation.
Their implementation will be tracked in #8027.
naga/src/valid/type.rs JJ: JJ: Lines starting with "JJ:" (like this
one) will be removed.
This adds HLSL backend support for `ImageClass::External` (ie WGSL's
`external_texture` texture type).
For each external texture global variable in the IR, we declare 3
`Texture2D` globals as well as a `cbuffer` for the params. The
additional bindings required by these are found in the newly added
`external_texture_binding_map`. Unique names for each can be obtained
using `NameKey::ExternalTextureGlobalVariable`.
For functions that contain ImageQuery::Size, ImageLoad, or ImageSample
expressions for external textures, ensure we have generated wrapper
functions for those expressions. When emitting code for the
expressions themselves, simply insert a call to the wrapper function.
For size queries, we return the value provided in the params
struct. If that value is [0, 0] then we query the size of the plane 0
texture and return that.
For load and sample, we sample the textures based on the number of
planes specified in the params struct. If there is more than one plane
we additionally perform YUV to RGB conversion using the provided
matrix.
Unfortunately HLSL does not allow structs to contain textures, meaning
we are unable to wrap the 3 textures and params struct variables in a
single variable that can be passed around.
For our wrapper functions we therefore ensure they take the three
textures and the params as consecutive arguments. Likewise, when
declaring user-defined functions with external texture arguments, we
expand the single external texture argument into 4 consecutive
arguments. (Using NameKey::ExternalTextureFunctionArgument to ensure
unique names for each.)
Thankfully external textures can only be used as either global
variables or function arguments. This means we only have to handle the
`Expression::GlobalVariable` and `Expression::FunctionArgument` cases
of `write_expr()`. Since in both cases we know the external texture
can only be an argument to either a user-defined function or one of
our wrapper functions, we can simply emit the names of the variables
for each three textures and the params struct in a comma-separated
list.
Adds new `NameKey` variants `ExternalTextureGlobalVariable` and
`ExternalTextureFunctionArgument`, like their non-external-texture
cousins but additionally keyed by either being a specific plane index
or params buffer.
For each external texture global variable or function argument reserve
additional names for 3 planes and the params buffer. For Naga backends
which must represent external textures as multiple variables/arguments,
this will allow them to uniquely name each one.
During wgsl lowering, if we encounter an external texture type then
generate the `ExternalTextureParams` struct. This will be required by
most Naga backends to implement external textures.
This type is not actually used by wgsl-in or the IR. However,
generating it in Naga IR ensures tricky details such as member
alignment are handled for us.
wgsl-out must ensure it does *not* generate code for this type, as it
handles external textures natively.
* [naga wgsl-in wgsl-out] WGSL support for texture_external texture type
Make wgsl-in correctly parse `texture_external` texture declarations,
and allow such textures to be used in `textureDimensions()`,
`textureSampleBaseClampToEdge()`, and `textureLoad()` function
calls. In IR these are represented by the `ImageClass::External` image
class, which is a 2D, non-multisampled, non-mipmapped, float-sampled
image.
Adds a new Capability `TEXTURE_EXTERNAL` and ensure validation rejects
shaders containing external textures if this capability flag is not
set. This capability is enabled for validation by wgpu devices which
support the `TEXTURE_EXTERNAL` feature (currently only when using the
noop backend), and by the Naga CLI when validating-only or when
outputting WGSL.
The WGSL backend can of course emit `ImageClass::External` images
directly as `texture_external` textures. Other backends are, for now,
unimplemented.
Lastly, we add a snapshot test covering all the valid uses of a
texture_external texture. These are:
- As a global variable declaration
- As an argument to the built-in functions `textureDimensions()`,
`textureSampleBaseClampToEdge()`, and `textureLoad()`
- As an argument to user-defined function declarations and calls.
We keep these in their own test so that we can control which targets
to run them against (currently WGSL and IR). When external textures
are supported by all Naga backends we can, if so inclined, integrate
these with existing texture tests.
* fixup! [naga wgsl-in wgsl-out] WGSL support for texture_external texture type
* fixup! [naga wgsl-in wgsl-out] WGSL support for texture_external texture type
---------
Co-authored-by: Jim Blandy <jimb@red-bean.com>
Adds a new flag to the IR indicating when image sample coordinates are
to be clamped. Adds wgsl-in support for parsing and lowering to
IR. Validation ensures this flag is only used when sampling a 2D
non-arrayed sampled texture, without offset, gather, or depth
comparison. This matches the WGSL requirements, with the exception of
supporting `texture_external` textures, which will follow in a later
patch.
SPIRV, HLSL, and Metal backends are supported so far, with GLSL left
for a follow up. (In GLSL the texture will simply be sampled without
the coordinates being clamped.)
It may seem unfortunate to have to handle this separately for each
backend, and indeed it would have been possible to implement this simply
in the WGSL frontend. However, future patches will add support for using
textureSampleBaseClampToEdge() with external textures, which will
actually have to be handled by each backend. This patch is laying the
groundwork for that.
In the SPIR-V front end, when generating Naga IR `Struct` types to
represent a SPIR-V entry point's `Output` variables, instead of saying
"0xFFFF, // shouldn't matter", follow the usual rules in assigning
struct member offsets and computing an overall
size (`TypeInner::Struct::span`) for the resulting struct type.
Adds a mode to compaction that removes unused functions, global
variables, and named types and overrides. This mode is used
everywhere except the compaction at the end of lowering, where
it is important to preserve unused items for type checking and
other validation of the module.
Pruning all but the active entry point and then compacting makes
`process_overrides` tolerant of missing values for overrides that are
not used by the active entry point.
Fixes#5885
* [wgsl-in,ir] add support for parsing rust-style doc comments
* rename relevant items to `doc_comments` (or variations of it)
* address comments
* remove `next_until`
* rename `save_doc_comments` to `ignore_doc_comments`
* expand snapshot test and ignore blankspace when accumulating doc comments
* make tokenizer more straightforward
---------
Co-authored-by: teoxoy <28601907+teoxoy@users.noreply.github.com>
Implements more direct conversions between 32-bit integers and 4x8-bit
integer vectors using bit casting to/from `packed_[u]char4` when on
MSL 2.1+ (older versions of MSL don't seem to support these bit casts).
- `unpack4x{I, U}8(x)` becomes `[u]int4(as_type<packed_[u]char4>(x))`;
- `pack4x{I, U}8(x)` becomes `as_type<uint>(packed_[u]char4(x))`; and
- `pack4x{I, U}8Clamp(x)` becomes
`as_type<uint>(packed_uchar4(metal::clamp(x, 0, 255)))`.
These bit casts match the WGSL spec for these functions because Metal
runs on little-endian machines.
Emits vectorized SPIR-V code for the WGSL functions `unpack4xI8`,
`unpack4xU8`, `pack4xI8`, `pack4xU8`, `pack4xI8Clamp`, and
`pack4xU8Clamp` if `Capability::Int8` is available.
Exploits the following facts about SPIR-V ops:
- `SClamp`, `UClamp`, and `OpUConvert` accept vector arguments, in which
case results are computed per component; and
- `OpBitcast` can cast between vectors and scalars, with a well-defined
bit order that matches that required by the WGSL spec, see below.
WGSL spec for `pack4xI8` [1]:
> Component e[i] of the input is mapped to bits 8 x i through 8 x i + 7
> of the result.
SPIR-V spec for `OpBitcast` [2]:
> Within this mapping, any single component of `S` [remark: the type
> with fewer but wider components] (mapping to multiple components of
> `L` [remark: the type with more but narrower components]) maps its
> lower-ordered bits to the lower-numbered components of `L`.
[1] https://www.w3.org/TR/WGSL/#pack4xI8-builtin
[2] https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitcast
* Potentially optimize `dot4{I,U}8Packed` on Metal
This might allow the Metal compiler to emit faster code (but that's not
confirmed). See
<https://github.com/gpuweb/gpuweb/issues/2677#issuecomment-1713292226>
for the optimization. The limitation to Metal 2.1+ is discussed here:
<https://github.com/gfx-rs/wgpu/pull/7574#issuecomment-2835464472>.
* [naga] Factor out new part of `put_block` on msl
CI on test failed because the latest changes to `put_block` made its
stack too big. Factoring out the new code into a separate method fixes
this issue.
