Implementing it is tracked by #6495. In the mean time it is okay for it
to be a no-op. `destroy` just recovers the resources earlier than garbage
collection might, doing nothing is not a leak.
* [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>
* Make wgpu-core's copy_buffer_to_buffer `size` parameter optional
* Make the copy size optional in more places
* Fix for webgpu backend
* [deno_webgpu] Support additional copyBufferToBuffer signatures
* Add changelog entry
* Add copyBufferToBuffer tests to CTS test list
(This doesn't actually enable the tests for the new overloads, because
of a different error reporting issue that affects many CTS tests
including these. But if you run the tests for the new overloads
manually, before and after the fix, you can see that the behavior has
changed.)
* Reproducible formula for vendoring modified webgpu-sys
Commit the updated vendor command in all the files for consistency.
Separates the Vulkan feature sets
`VkPhysicalDeviceShaderFloat16Int8Features` and
`VkPhysicalDevice16BitStorageFeatures`, which previously were used
"together, or not at all".
This commit should not change any behavior yet, but I'd like to run full
CI tests on it for now. If the CI tests pass, I'll use this separation
to enable the `shader_int8` feature separately from the rest of the
features to enable optimizations of `[un]pack4x{I,U}8[Clamp]` on SPIR-V.
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
