nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-21 22:33:49 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'trunk' into objc2-metal

Browse Source
This commit is contained in:
Mads Marquart 2024-10-04 15:45:48 +02:00
commit 1e04db3d27
182 changed files with 3972 additions and 1799 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
.github/CODEOWNERS vendored
View File

@ -1,10 +1,17 @@
* @gfx-rs/wgpu
/cts_runner/ @gfx-rs/deno @gfx-rs/wgpu
/deno_webgpu/ @gfx-rs/deno @gfx-rs/wgpu
/cts_runner/ @crowlkats @gfx-rs/wgpu
/deno_webgpu/ @crowlkats @gfx-rs/wgpu
/naga/ @gfx-rs/naga
/naga-cli/ @gfx-rs/naga
# Both wgpu and naga teams are owners of naga infrastructure so
# either team can review changes to deps and docs.
naga/Cargo.toml @gfx-rs/wgpu @gfx-rs/naga
naga/README.md @gfx-rs/wgpu @gfx-rs/naga
naga/CHANGELOG.md @gfx-rs/wgpu @gfx-rs/naga
naga-cli/Cargo.toml @gfx-rs/wgpu @gfx-rs/naga
# We leave the codeowners empty for the changelog, so naga changes
# don't trigger wgpu reviews and vise versa.
/CHANGELOG.md

21
.github/workflows/ci.yml vendored
View File

@ -13,21 +13,23 @@ env:
#
# Sourced from https://vulkan.lunarg.com/sdk/home#linux
VULKAN_SDK_VERSION: "1.3.268"
#
# We don't include the 4th version number, as it's not used in any URL.
VULKAN_SDK_VERSION: "1.3.290"
# Sourced from https://www.nuget.org/packages/Microsoft.Direct3D.WARP
WARP_VERSION: "1.0.8"
WARP_VERSION: "1.0.13"
# Sourced from https://github.com/microsoft/DirectXShaderCompiler/releases
#
# Must also be changed in shaders.yaml
DXC_RELEASE: "v1.7.2308"
DXC_FILENAME: "dxc_2023_08_14.zip"
DXC_RELEASE: "v1.8.2407"
DXC_FILENAME: "dxc_2024_07_31_clang_cl.zip"
# Sourced from https://archive.mesa3d.org/. Bumping this requires
# updating the mesa build in https://github.com/gfx-rs/ci-build and creating a new release.
MESA_VERSION: "23.3.1"
MESA_VERSION: "24.2.3"
# Corresponds to https://github.com/gfx-rs/ci-build/releases
CI_BINARY_BUILD: "build18"
CI_BINARY_BUILD: "build19"
# We sometimes need nightly to use special things in CI.
#
@ -311,7 +313,6 @@ jobs:
rustup override set ${{ env.CORE_MSRV }}
cargo -V
# Use special toolchain for rustdoc, see https://github.com/gfx-rs/wgpu/issues/4905
- name: Install Nightly Toolchain
run: |
rustup toolchain install ${{ env.NIGHTLY_VERSION }} --no-self-update --profile=minimal --component clippy
@ -442,7 +443,7 @@ jobs:
dxc --version
curl.exe -L --retry 5 https://www.nuget.org/api/v2/package/Microsoft.Direct3D.WARP/$WARP_VERSION -o warp.zip
7z.exe e warp.zip -owarp build/native/amd64/d3d10warp.dll
7z.exe e warp.zip -owarp build/native/bin/x64/d3d10warp.dll
mkdir -p target/llvm-cov-target/debug/deps
@ -554,7 +555,7 @@ jobs:
cargo llvm-cov report --lcov --output-path lcov.info
- name: upload coverage report to codecov
uses: codecov/codecov-action@v3
uses: codecov/codecov-action@v4
if: steps.coverage.outcome == 'success'
with:
files: lcov.info
@ -619,7 +620,7 @@ jobs:
cargo fmt --manifest-path xtask/Cargo.toml -- --check
- name: Check for typos
uses: crate-ci/typos@v1.24.5
uses: crate-ci/typos@v1.24.6
check-cts-runner:
# runtime is normally 2 minutes

2
.github/workflows/docs.yml vendored
View File

@ -41,7 +41,7 @@ jobs:
if: ${{ failure() }}
- name: Deploy the docs
uses: JamesIves/github-pages-deploy-action@v4.6.4
uses: JamesIves/github-pages-deploy-action@v4.6.8
if: github.ref == 'refs/heads/trunk'
with:
token: ${{ secrets.WEB_DEPLOY }}

2
.github/workflows/publish.yml vendored
View File

@ -41,7 +41,7 @@ jobs:
run: cargo xtask run-wasm --no-serve
- name: Deploy WebGPU examples
uses: JamesIves/github-pages-deploy-action@v4.6.4
uses: JamesIves/github-pages-deploy-action@v4.6.8
if: github.ref == 'refs/heads/trunk'
with:
token: ${{ secrets.WEB_DEPLOY }}

10
.github/workflows/shaders.yml vendored
View File

@ -9,13 +9,17 @@ on:
env:
# Sourced from https://vulkan.lunarg.com/sdk/home#linux
VULKAN_SDK_VERSION: "1.3.268"
#
# We don't include the 4th version number, as it's not used in any URL.
#
# Held back from 1.3.290 by https://github.com/gfx-rs/wgpu/issues/6307
VULKAN_SDK_VERSION: "1.3.283"
# Sourced from https://github.com/microsoft/DirectXShaderCompiler/releases
#
# Must also be changed in ci.yaml
DXC_RELEASE: "v1.7.2308"
DXC_FILENAME: "dxc_2023_08_14.zip"
DXC_RELEASE: "v1.8.2407"
DXC_FILENAME: "dxc_2024_07_31_clang_cl.zip"
jobs:
naga-validate-windows:

39
CHANGELOG.md
View File

@ -27,11 +27,12 @@ Top level categories:
Bottom level categories:
- Naga
- General
- DX12
- Vulkan
- Metal
- GLES
- GLES / OpenGL
- WebGPU
- Emscripten
- Hal
@ -80,6 +81,13 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
- Support constant evaluation for `firstLeadingBit` and `firstTrailingBit` numeric built-ins in WGSL. Front-ends that translate to these built-ins also benefit from constant evaluation. By @ErichDonGubler in [#5101](https://github.com/gfx-rs/wgpu/pull/5101).
- Add `first` and `either` sampling types for `@interpolate(flat, …)` in WGSL. By @ErichDonGubler in [#6181](https://github.com/gfx-rs/wgpu/pull/6181).
- Support for more atomic ops in the SPIR-V frontend. By @schell in [#5824](https://github.com/gfx-rs/wgpu/pull/5824).
- Support local `const` declarations in WGSL. By @sagudev in [#6156](https://github.com/gfx-rs/wgpu/pull/6156).
- Implemented `const_assert` in WGSL. By @sagudev in [#6198](https://github.com/gfx-rs/wgpu/pull/6198).
#### General
- Add `VideoFrame` to `ExternalImageSource` enum. By @jprochazk in [#6170](https://github.com/gfx-rs/wgpu/pull/6170)
#### Vulkan
@ -92,6 +100,13 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
#### Naga
- Accept only `vec3` (not `vecN`) for the `cross` built-in. By @ErichDonGubler in [#6171](https://github.com/gfx-rs/wgpu/pull/6171).
- Configure `SourceLanguage` when enabling debug info in SPV-out. By @kvark in [#6256](https://github.com/gfx-rs/wgpu/pull/6256).
- Per-polygon and flat inputs should not be considered subgroup uniform. By @magcius in [#6276](https://github.com/gfx-rs/wgpu/pull/6276).
- Validate all swizzle components are either color (rgba) or dimension (xyzw) in WGSL. By @sagudev in [#6187](https://github.com/gfx-rs/wgpu/pull/6187).
- Fix detection of shl overflows to detect arithmetic overflows. By @sagudev in [#6186](https://github.com/gfx-rs/wgpu/pull/6186).
- Fix type parameters to vec/mat type constructors to also support aliases. By @sagudev in [#6189](https://github.com/gfx-rs/wgpu/pull/6189).
- Accept global `var`s without explicit type. By @sagudev in [#6199](https://github.com/gfx-rs/wgpu/pull/6199).
- Fix handling of phony statements, so they are actually emitted. By @sagudev in [#6328](https://github.com/gfx-rs/wgpu/pull/6328).
#### General
@ -104,12 +119,23 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
- Deduplicate bind group layouts that are created from pipelines with "auto" layouts. By @teoxoy [#6049](https://github.com/gfx-rs/wgpu/pull/6049)
- Fix crash when dropping the surface after the device. By @wumpf in [#6052](https://github.com/gfx-rs/wgpu/pull/6052)
- Fix error message that is thrown in create_render_pass to no longer say `compute_pass`. By @matthew-wong1 [#6041](https://github.com/gfx-rs/wgpu/pull/6041)
- Add `VideoFrame` to `ExternalImageSource` enum. By @jprochazk in [#6170](https://github.com/gfx-rs/wgpu/pull/6170)
- Document `wgpu_hal` bounds-checking promises, and adapt `wgpu_core`'s lazy initialization logic to the slightly weaker-than-expected guarantees. By @jimblandy in [#6201](https://github.com/gfx-rs/wgpu/pull/6201)
- Raise validation error instead of panicking in `{Render,Compute}Pipeline::get_bind_group_layout` on native / WebGL. By @bgr360 in [#6280](https://github.com/gfx-rs/wgpu/pull/6280).
- **BREAKING**: Remove the last exposed C symbols in project, located in `wgpu_core::render::bundle::bundle_ffi`, to allow multiple versions of WGPU to compile together. By @ErichDonGubler in [#6272](https://github.com/gfx-rs/wgpu/pull/6272).
#### GLES / OpenGL
- Fix GL debug message callbacks not being properly cleaned up (causing UB). By @Imberflur in [#6114](https://github.com/gfx-rs/wgpu/pull/6114)
- Fix calling `slice::from_raw_parts` with unaligned pointers in push constant handling. By @Imberflur in [#6341](https://github.com/gfx-rs/wgpu/pull/6341)
#### WebGPU
- Fix JS `TypeError` exception in `Instance::request_adapter` when browser doesn't support WebGPU but `wgpu` not compiled with `webgl` support. By @bgr360 in [#6197](https://github.com/gfx-rs/wgpu/pull/6197).
#### Vulkan
- Vulkan debug labels assumed no interior nul byte. By @DJMcNab in [#6257](https://github.com/gfx-rs/wgpu/pull/6257)
- Add `.index_type(vk::IndexType::NONE_KHR)` when creating `AccelerationStructureGeometryTrianglesDataKHR` in the raytraced triangle example to prevent a validation error. By @Vecvec in [#6282](https://github.com/gfx-rs/wgpu/pull/6282)
#### Metal
@ -125,6 +151,11 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
- Change the inconsistent `DropGuard` based API on Vulkan and GLES to a consistent, callback-based one. By @jerzywilczek in [#6164](https://github.com/gfx-rs/wgpu/pull/6164)
### Documentation
- Removed some OpenGL and Vulkan references from `wgpu-types` documentation. Fixed Storage texel types in examples. By @Nelarius in [#6271](https://github.com/gfx-rs/wgpu/pull/6271)
- Used `wgpu::include_wgsl!(…)` more in examples and tests. By @ErichDonGubler in [#6326](https://github.com/gfx-rs/wgpu/pull/6326).
### Dependency Updates
#### GLES
@ -137,6 +168,10 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
- Replace `winapi` code to use the `windows` crate. By @MarijnS95 in [#5956](https://github.com/gfx-rs/wgpu/pull/5956) and [#6173](https://github.com/gfx-rs/wgpu/pull/6173)
#### HAL
- Update `parking_lot` to `0.12`. By @mahkoh in [#6287](https://github.com/gfx-rs/wgpu/pull/6287)
## 22.0.0 (2024-07-17)
### Overview

164
Cargo.lock generated
View File

@ -160,9 +160,9 @@ dependencies = [
[[package]]
name = "anyhow"
version = "1.0.87"
version = "1.0.89"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "10f00e1f6e58a40e807377c75c6a7f97bf9044fab57816f2414e6f5f4499d7b8"
checksum = "86fdf8605db99b54d3cd748a44c6d04df638eb5dafb219b135d0149bd0db01f6"
[[package]]
name = "arbitrary"
@ -206,9 +206,9 @@ dependencies = [
[[package]]
name = "arrayref"
version = "0.3.8"
version = "0.3.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "9d151e35f61089500b617991b791fc8bfd237ae50cd5950803758a179b41e67a"
checksum = "76a2e8124351fda1ef8aaaa3bbd7ebbcb486bbcd4225aca0aa0d84bb2db8fecb"
[[package]]
name = "arrayvec"
@ -236,9 +236,9 @@ dependencies = [
[[package]]
name = "async-trait"
version = "0.1.82"
version = "0.1.83"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a27b8a3a6e1a44fa4c8baf1f653e4172e81486d4941f2237e20dc2d0cf4ddff1"
checksum = "721cae7de5c34fbb2acd27e21e6d2cf7b886dce0c27388d46c4e6c47ea4318dd"
dependencies = [
"proc-macro2",
"quote",
@ -253,9 +253,9 @@ checksum = "1505bd5d3d116872e7271a6d4e16d81d0c8570876c8de68093a09ac269d8aac0"
[[package]]
name = "autocfg"
version = "1.3.0"
version = "1.4.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0c4b4d0bd25bd0b74681c0ad21497610ce1b7c91b1022cd21c80c6fbdd9476b0"
checksum = "ace50bade8e6234aa140d9a2f552bbee1db4d353f69b8217bc503490fc1a9f26"
[[package]]
name = "backtrace"
@ -399,9 +399,9 @@ checksum = "1fd0f2584146f6f2ef48085050886acf353beff7305ebd1ae69500e27c67f64b"
[[package]]
name = "bytes"
version = "1.7.1"
version = "1.7.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8318a53db07bb3f8dca91a600466bdb3f2eaadeedfdbcf02e1accbad9271ba50"
checksum = "428d9aa8fbc0670b7b8d6030a7fadd0f86151cae55e4dbbece15f3780a3dfaf3"
[[package]]
name = "calloop"
@ -437,9 +437,9 @@ checksum = "37b2a672a2cb129a2e41c10b1224bb368f9f37a2b16b612598138befd7b37eb5"
[[package]]
name = "cc"
version = "1.1.18"
version = "1.1.22"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b62ac837cdb5cb22e10a256099b4fc502b1dfe560cb282963a974d7abd80e476"
checksum = "9540e661f81799159abee814118cc139a2004b3a3aa3ea37724a1b66530b90e0"
dependencies = [
"jobserver",
"libc",
@ -502,9 +502,9 @@ dependencies = [
[[package]]
name = "clap"
version = "4.5.17"
version = "4.5.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3e5a21b8495e732f1b3c364c9949b201ca7bae518c502c80256c96ad79eaf6ac"
checksum = "b0956a43b323ac1afaffc053ed5c4b7c1f1800bacd1683c353aabbb752515dd3"
dependencies = [
"clap_builder",
"clap_derive",
@ -512,9 +512,9 @@ dependencies = [
[[package]]
name = "clap_builder"
version = "4.5.17"
version = "4.5.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8cf2dd12af7a047ad9d6da2b6b249759a22a7abc0f474c1dae1777afa4b21a73"
checksum = "4d72166dd41634086d5803a47eb71ae740e61d84709c36f3c34110173db3961b"
dependencies = [
"anstream",
"anstyle",
@ -524,9 +524,9 @@ dependencies = [
[[package]]
name = "clap_derive"
version = "4.5.13"
version = "4.5.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "501d359d5f3dcaf6ecdeee48833ae73ec6e42723a1e52419c79abf9507eec0a0"
checksum = "4ac6a0c7b1a9e9a5186361f67dfa1b88213572f427fb9ab038efb2bd8c582dab"
dependencies = [
"heck 0.5.0",
"proc-macro2",
@ -603,9 +603,9 @@ dependencies = [
[[package]]
name = "const_panic"
version = "0.2.9"
version = "0.2.10"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7782af8f90fe69a4bb41e460abe1727d493403d8b2cc43201a3a3e906b24379f"
checksum = "013b6c2c3a14d678f38cd23994b02da3a1a1b6a5d1eedddfe63a5a5f11b13a81"
[[package]]
name = "convert_case"
@ -1058,9 +1058,9 @@ checksum = "e8c02a5121d4ea3eb16a80748c74f5549a5665e4c21333c6098f283870fbdea6"
[[package]]
name = "fdeflate"
version = "0.3.4"
version = "0.3.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "4f9bfee30e4dedf0ab8b422f03af778d9612b63f502710fc500a334ebe2de645"
checksum = "d8090f921a24b04994d9929e204f50b498a33ea6ba559ffaa05e04f7ee7fb5ab"
dependencies = [
"simd-adler32",
]
@ -1082,9 +1082,9 @@ checksum = "0ce7134b9999ecaf8bcd65542e436736ef32ddca1b3e06094cb6ec5755203b80"
[[package]]
name = "flate2"
version = "1.0.33"
version = "1.0.34"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "324a1be68054ef05ad64b861cc9eaf1d623d2d8cb25b4bf2cb9cdd902b4bf253"
checksum = "a1b589b4dc103969ad3cf85c950899926ec64300a1a46d76c03a6072957036f0"
dependencies = [
"crc32fast",
"miniz_oxide 0.8.0",
@ -1311,9 +1311,9 @@ checksum = "779ae4bf7e8421cf91c0b3b64e7e8b40b862fba4d393f59150042de7c4965a94"
[[package]]
name = "glow"
version = "0.14.0"
version = "0.14.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f865cbd94bd355b89611211e49508da98a1fce0ad755c1e8448fb96711b24528"
checksum = "2f4a888dbe8181a7535853469c21c67ca9a1cea9460b16808fc018ea9e55d248"
dependencies = [
"js-sys",
"slotmap",
@ -1675,9 +1675,9 @@ checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe"
[[package]]
name = "libc"
version = "0.2.158"
version = "0.2.159"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d8adc4bb1803a324070e64a98ae98f38934d91957a99cfb3a43dcbc01bc56439"
checksum = "561d97a539a36e26a9a5fad1ea11a3039a67714694aaa379433e580854bc3dc5"
[[package]]
name = "libfuzzer-sys"
@ -1789,9 +1789,9 @@ checksum = "78ca9ab1a0babb1e7d5695e3530886289c18cf2f87ec19a575a0abdce112e3a3"
[[package]]
name = "memmap2"
version = "0.9.4"
version = "0.9.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "fe751422e4a8caa417e13c3ea66452215d7d63e19e604f4980461212f3ae1322"
checksum = "fd3f7eed9d3848f8b98834af67102b720745c4ec028fcd0aa0239277e7de374f"
dependencies = [
"libc",
]
@ -1822,7 +1822,6 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b8a240ddb74feaf34a79a7add65a741f3167852fba007066dcac1ca548d89c08"
dependencies = [
"adler",
"simd-adler32",
]
[[package]]
@ -1832,6 +1831,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e2d80299ef12ff69b16a84bb182e3b9df68b5a91574d3d4fa6e41b65deec4df1"
dependencies = [
"adler2",
"simd-adler32",
]
[[package]]
@ -2141,9 +2141,12 @@ dependencies = [
[[package]]
name = "once_cell"
version = "1.19.0"
version = "1.20.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3fdb12b2476b595f9358c5161aa467c2438859caa136dec86c26fdd2efe17b92"
checksum = "82881c4be219ab5faaf2ad5e5e5ecdff8c66bd7402ca3160975c93b24961afd1"
dependencies = [
"portable-atomic",
]
[[package]]
name = "oorandom"
@ -2275,9 +2278,9 @@ checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184"
[[package]]
name = "pkg-config"
version = "0.3.30"
version = "0.3.31"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d231b230927b5e4ad203db57bbcbee2802f6bce620b1e4a9024a07d94e2907ec"
checksum = "953ec861398dccce10c670dfeaf3ec4911ca479e9c02154b3a215178c5f566f2"
[[package]]
name = "player"
@ -2323,15 +2326,15 @@ dependencies = [
[[package]]
name = "png"
version = "0.17.13"
version = "0.17.14"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "06e4b0d3d1312775e782c86c91a111aa1f910cbb65e1337f9975b5f9a554b5e1"
checksum = "52f9d46a34a05a6a57566bc2bfae066ef07585a6e3fa30fbbdff5936380623f0"
dependencies = [
"bitflags 1.3.2",
"crc32fast",
"fdeflate",
"flate2",
"miniz_oxide 0.7.4",
"miniz_oxide 0.8.0",
]
[[package]]
@ -2355,6 +2358,12 @@ version = "0.3.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "22686f4785f02a4fcc856d3b3bb19bf6c8160d103f7a99cc258bddd0251dc7f2"
[[package]]
name = "portable-atomic"
version = "1.9.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cc9c68a3f6da06753e9335d63e27f6b9754dd1920d941135b7ea8224f141adb2"
[[package]]
name = "pp-rs"
version = "0.2.1"
@ -2419,9 +2428,9 @@ checksum = "43d84d1d7a6ac92673717f9f6d1518374ef257669c24ebc5ac25d5033828be58"
[[package]]
name = "quick-xml"
version = "0.36.1"
version = "0.36.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "96a05e2e8efddfa51a84ca47cec303fac86c8541b686d37cac5efc0e094417bc"
checksum = "f7649a7b4df05aed9ea7ec6f628c67c9953a43869b8bc50929569b2999d443fe"
dependencies = [
"memchr",
]
@ -2526,14 +2535,14 @@ dependencies = [
[[package]]
name = "regex"
version = "1.10.6"
version = "1.11.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "4219d74c6b67a3654a9fbebc4b419e22126d13d2f3c4a07ee0cb61ff79a79619"
checksum = "38200e5ee88914975b69f657f0801b6f6dccafd44fd9326302a4aaeecfacb1d8"
dependencies = [
"aho-corasick",
"memchr",
"regex-automata 0.4.6",
"regex-syntax 0.8.3",
"regex-automata 0.4.8",
"regex-syntax 0.8.5",
]
[[package]]
@ -2547,13 +2556,13 @@ dependencies = [
[[package]]
name = "regex-automata"
version = "0.4.6"
version = "0.4.8"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "86b83b8b9847f9bf95ef68afb0b8e6cdb80f498442f5179a29fad448fcc1eaea"
checksum = "368758f23274712b504848e9d5a6f010445cc8b87a7cdb4d7cbee666c1288da3"
dependencies = [
"aho-corasick",
"memchr",
"regex-syntax 0.8.3",
"regex-syntax 0.8.5",
]
[[package]]
@ -2564,9 +2573,9 @@ checksum = "f162c6dd7b008981e4d40210aca20b4bd0f9b60ca9271061b07f78537722f2e1"
[[package]]
name = "regex-syntax"
version = "0.8.3"
version = "0.8.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "adad44e29e4c806119491a7f06f03de4d1af22c3a680dd47f1e6e179439d1f56"
checksum = "2b15c43186be67a4fd63bee50d0303afffcef381492ebe2c5d87f324e1b8815c"
[[package]]
name = "renderdoc-sys"
@ -2627,9 +2636,9 @@ dependencies = [
[[package]]
name = "rustix"
version = "0.38.36"
version = "0.38.37"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3f55e80d50763938498dd5ebb18647174e0c76dc38c5505294bb224624f30f36"
checksum = "8acb788b847c24f28525660c4d7758620a7210875711f79e7f663cc152726811"
dependencies = [
"bitflags 2.6.0",
"errno",
@ -2945,9 +2954,9 @@ dependencies = [
[[package]]
name = "syn"
version = "2.0.77"
version = "2.0.79"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "9f35bcdf61fd8e7be6caf75f429fdca8beb3ed76584befb503b1569faee373ed"
checksum = "89132cd0bf050864e1d38dc3bbc07a0eb8e7530af26344d3d2bbbef83499f590"
dependencies = [
"proc-macro2",
"quote",
@ -2965,18 +2974,18 @@ dependencies = [
[[package]]
name = "thiserror"
version = "1.0.63"
version = "1.0.64"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c0342370b38b6a11b6cc11d6a805569958d54cfa061a29969c3b5ce2ea405724"
checksum = "d50af8abc119fb8bb6dbabcfa89656f46f84aa0ac7688088608076ad2b459a84"
dependencies = [
"thiserror-impl",
]
[[package]]
name = "thiserror-impl"
version = "1.0.63"
version = "1.0.64"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a4558b58466b9ad7ca0f102865eccc95938dca1a74a856f2b57b6629050da261"
checksum = "08904e7672f5eb876eaaf87e0ce17857500934f4981c4a0ab2b4aa98baac7fc3"
dependencies = [
"proc-macro2",
"quote",
@ -3099,9 +3108,9 @@ checksum = "0dd7358ecb8fc2f8d014bf86f6f638ce72ba252a2c3a2572f2a795f1d23efb41"
[[package]]
name = "toml_edit"
version = "0.22.20"
version = "0.22.22"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "583c44c02ad26b0c3f3066fe629275e50627026c51ac2e595cca4c230ce1ce1d"
checksum = "4ae48d6208a266e853d946088ed816055e556cc6028c5e8e2b84d9fa5dd7c7f5"
dependencies = [
"indexmap",
"toml_datetime",
@ -3159,9 +3168,9 @@ dependencies = [
[[package]]
name = "tracy-client"
version = "0.17.3"
version = "0.17.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "373db47331c3407b343538df77eea2516884a0b126cdfb4b135acfd400015dd7"
checksum = "746b078c6a09ebfd5594609049e07116735c304671eaab06ce749854d23435bc"
dependencies = [
"loom",
"once_cell",
@ -3170,9 +3179,9 @@ dependencies = [
[[package]]
name = "tracy-client-sys"
version = "0.24.0"
version = "0.24.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "49cf0064dcb31c99aa1244c1b93439359e53f72ed217eef5db50abd442241e9a"
checksum = "68613466112302fdbeabc5fa55f7d57462a0b247d5a6b7d7e09401fb471a144d"
dependencies = [
"cc",
]
@ -3232,42 +3241,42 @@ checksum = "08f95100a766bf4f8f28f90d77e0a5461bbdb219042e7679bebe79004fed8d75"
[[package]]
name = "unicode-id-start"
version = "1.2.0"
version = "1.3.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "bc3882f69607a2ac8cc4de3ee7993d8f68bb06f2974271195065b3bd07f2edea"
checksum = "97e2a3c5fc9de285c0e805d98eba666adb4b2d9e1049ce44821ff7707cc34e91"
[[package]]
name = "unicode-ident"
version = "1.0.12"
version = "1.0.13"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3354b9ac3fae1ff6755cb6db53683adb661634f67557942dea4facebec0fee4b"
checksum = "e91b56cd4cadaeb79bbf1a5645f6b4f8dc5bde8834ad5894a8db35fda9efa1fe"
[[package]]
name = "unicode-normalization"
version = "0.1.23"
version = "0.1.24"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a56d1686db2308d901306f92a263857ef59ea39678a5458e7cb17f01415101f5"
checksum = "5033c97c4262335cded6d6fc3e5c18ab755e1a3dc96376350f3d8e9f009ad956"
dependencies = [
"tinyvec",
]
[[package]]
name = "unicode-segmentation"
version = "1.11.0"
version = "1.12.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d4c87d22b6e3f4a18d4d40ef354e97c90fcb14dd91d7dc0aa9d8a1172ebf7202"
checksum = "f6ccf251212114b54433ec949fd6a7841275f9ada20dddd2f29e9ceea4501493"
[[package]]
name = "unicode-width"
version = "0.1.13"
version = "0.1.14"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0336d538f7abc86d282a4189614dfaa90810dfc2c6f6427eaf88e16311dd225d"
checksum = "7dd6e30e90baa6f72411720665d41d89b9a3d039dc45b8faea1ddd07f617f6af"
[[package]]
name = "unicode-xid"
version = "0.2.5"
version = "0.2.6"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "229730647fbc343e3a80e463c1db7f78f3855d3f3739bee0dda773c9a037c90a"
checksum = "ebc1c04c71510c7f702b52b7c350734c9ff1295c464a03335b00bb84fc54f853"
[[package]]
name = "url"
@ -3689,6 +3698,7 @@ dependencies = [
"bit-set",
"bitflags 2.6.0",
"block2 0.5.1",
"bytemuck",
"cfg-if",
"cfg_aliases",
"env_logger",
@ -4170,9 +4180,9 @@ dependencies = [
[[package]]
name = "winnow"
version = "0.6.18"
version = "0.6.20"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "68a9bda4691f099d435ad181000724da8e5899daa10713c2d432552b9ccd3a6f"
checksum = "36c1fec1a2bb5866f07c25f68c26e565c4c200aebb96d7e55710c19d3e8ac49b"
dependencies = [
"memchr",
]

14
Cargo.toml
View File

@ -70,13 +70,13 @@ path = "./naga"
version = "22.0.0"
[workspace.dependencies]
anyhow = "1.0.87"
anyhow = "1.0.89"
argh = "0.1.5"
arrayvec = "0.7"
bincode = "1"
bit-vec = "0.8"
bitflags = "2.6"
bytemuck = { version = "1.18", features = ["derive"] }
bytemuck = { version = "1.18" }
cfg_aliases = "0.1"
cfg-if = "1"
criterion = "0.5"
@ -104,14 +104,14 @@ nanorand = { version = "0.7", default-features = false, features = ["wyrand"] }
noise = { version = "0.8", git = "https://github.com/Razaekel/noise-rs.git", rev = "c6942d4fb70af26db4441edcf41f90fa115333f2" }
nv-flip = "0.1"
obj = "0.10"
once_cell = "1.19.0"
parking_lot = ">=0.11, <0.13" # parking_lot 0.12 switches from `winapi` to `windows`; permit either
once_cell = "1.20.1"
parking_lot = "0.12.1"
pico-args = { version = "0.5.0", features = [
"eq-separator",
"short-space-opt",
"combined-flags",
] }
png = "0.17.11"
png = "0.17.14"
pollster = "0.3"
profiling = { version = "1", default-features = false }
raw-window-handle = "0.6"
@ -125,7 +125,7 @@ smallvec = "1"
static_assertions = "1.1.0"
strum = { version = "0.25.0", features = ["derive"] }
tracy-client = "0.17"
thiserror = "1.0.63"
thiserror = "1.0.64"
wgpu = { version = "22.0.0", path = "./wgpu", default-features = false }
wgpu-core = { version = "22.0.0", path = "./wgpu-core" }
wgpu-macros = { version = "22.0.0", path = "./wgpu-macros" }
@ -194,7 +194,7 @@ windows-core = { version = "0.58", default-features = false }
# Gles dependencies
khronos-egl = "6"
glow = "0.14.0"
glow = "0.14.1"
glutin = { version = "0.31", default-features = false }
glutin-winit = { version = "0.4", default-features = false }
glutin_wgl_sys = "0.6"

2
benches/benches/root.rs
View File

@ -39,7 +39,7 @@ impl DeviceState {
let adapter_info = adapter.get_info();
eprintln!("{:?}", adapter_info);
eprintln!("{adapter_info:?}");
let (device, queue) = block_on(adapter.request_device(
&wgpu::DeviceDescriptor {

2
benches/benches/shader.rs
View File

@ -40,7 +40,7 @@ impl Inputs {
_ => continue,
},
Err(e) => {
eprintln!("Skipping file: {:?}", e);
eprintln!("Skipping file: {e:?}");
continue;
}
}

5
deno_webgpu/lib.rs
View File

@ -401,10 +401,7 @@ pub fn op_webgpu_request_adapter(
force_fallback_adapter,
compatible_surface: None, // windowless
};
let res = instance.request_adapter(
&descriptor,
wgpu_core::instance::AdapterInputs::Mask(backends, |_| None),
);
let res = instance.request_adapter(&descriptor, backends, None);
let adapter = match res {
Ok(adapter) => adapter,

2
deno_webgpu/queue.rs
View File

@ -44,7 +44,7 @@ pub fn op_webgpu_queue_submit(
})
.collect::<Result<Vec<_>, AnyError>>()?;
let maybe_err = instance.queue_submit(queue, &ids).err();
let maybe_err = instance.queue_submit(queue, &ids).err().map(|(_idx, e)| e);
for rid in command_buffers {
let resource = state.resource_table.take::<WebGpuCommandBuffer>(rid)?;

2
examples/Cargo.toml
View File

@ -29,7 +29,7 @@ webgl = ["wgpu/webgl"]
webgpu = ["wgpu/webgpu"]
[dependencies]
bytemuck.workspace = true
bytemuck = { workspace = true, features = ["derive"] }
cfg-if.workspace = true
encase = { workspace = true, features = ["glam"] }
flume.workspace = true

12
examples/src/boids/mod.rs
View File

@ -2,7 +2,7 @@
// adapted from https://github.com/austinEng/webgpu-samples/blob/master/src/examples/computeBoids.ts
use nanorand::{Rng, WyRand};
use std::{borrow::Cow, mem::size_of};
use std::mem::size_of;
use wgpu::util::DeviceExt;
// number of boid particles to simulate
@ -43,14 +43,8 @@ impl crate::framework::Example for Example {
device: &wgpu::Device,
_queue: &wgpu::Queue,
) -> Self {
let compute_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("compute.wgsl"))),
});
let draw_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("draw.wgsl"))),
});
let compute_shader = device.create_shader_module(wgpu::include_wgsl!("compute.wgsl"));
let draw_shader = device.create_shader_module(wgpu::include_wgsl!("draw.wgsl"));
// buffer for simulation parameters uniform

15
examples/src/conservative_raster/mod.rs
View File

@ -1,5 +1,3 @@
use std::borrow::Cow;
const RENDER_TARGET_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba8UnormSrgb;
struct Example {
@ -83,12 +81,8 @@ impl crate::framework::Example for Example {
push_constant_ranges: &[],
});
let shader_triangle_and_lines = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!(
"triangle_and_lines.wgsl"
))),
});
let shader_triangle_and_lines =
device.create_shader_module(wgpu::include_wgsl!("triangle_and_lines.wgsl"));
let pipeline_triangle_conservative =
device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
@ -203,10 +197,7 @@ impl crate::framework::Example for Example {
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("upscale.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("upscale.wgsl"));
(
device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("Upscale"),

7
examples/src/cube/mod.rs
View File

@ -1,5 +1,5 @@
use bytemuck::{Pod, Zeroable};
use std::{borrow::Cow, f32::consts, mem::size_of};
use std::{f32::consts, mem::size_of};
use wgpu::util::DeviceExt;
#[repr(C)]
@ -216,10 +216,7 @@ impl crate::framework::Example for Example {
label: None,
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let vertex_buffers = [wgpu::VertexBufferLayout {
array_stride: vertex_size as wgpu::BufferAddress,

7
examples/src/hello_compute/mod.rs
View File

@ -1,4 +1,4 @@
use std::{borrow::Cow, mem::size_of_val, str::FromStr};
use std::{mem::size_of_val, str::FromStr};
use wgpu::util::DeviceExt;
// Indicates a u32 overflow in an intermediate Collatz value
@ -66,10 +66,7 @@ async fn execute_gpu_inner(
numbers: &[u32],
) -> Option<Vec<u32>> {
// Loads the shader from WGSL
let cs_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let cs_module = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
// Gets the size in bytes of the buffer.
let size = size_of_val(numbers) as wgpu::BufferAddress;

5
examples/src/hello_synchronization/mod.rs
View File

@ -56,10 +56,7 @@ async fn execute(
let mut local_patient_workgroup_results = vec![0u32; result_vec_size];
let mut local_hasty_workgroup_results = local_patient_workgroup_results.clone();
let shaders_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!("shaders.wgsl"))),
});
let shaders_module = device.create_shader_module(wgpu::include_wgsl!("shaders.wgsl"));
let storage_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: None,

5
examples/src/hello_workgroups/mod.rs
View File

@ -41,10 +41,7 @@ async fn run() {
.await
.unwrap();
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let storage_buffer_a = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: None,

2
examples/src/main.rs
View File

@ -201,7 +201,7 @@ fn print_unknown_example(_result: Option<String>) {}
#[cfg(not(target_arch = "wasm32"))]
fn print_unknown_example(result: Option<String>) {
if let Some(example) = result {
println!("Unknown example: {}", example);
println!("Unknown example: {example}");
} else {
println!("Please specify an example as the first argument!");
}

12
examples/src/mipmap/mod.rs
View File

@ -1,5 +1,5 @@
use bytemuck::{Pod, Zeroable};
use std::{borrow::Cow, f32::consts, mem::size_of};
use std::{f32::consts, mem::size_of};
use wgpu::util::DeviceExt;
const TEXTURE_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba8UnormSrgb;
@ -81,10 +81,7 @@ impl Example {
query_sets: &Option<QuerySets>,
mip_count: u32,
) {
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("blit.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("blit.wgsl"));
let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("blit"),
@ -281,10 +278,7 @@ impl crate::framework::Example for Example {
});
// Create the render pipeline
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("draw.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("draw.wgsl"));
let draw_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("draw"),

7
examples/src/msaa_line/mod.rs
View File

@ -7,7 +7,7 @@
//! * Set the primitive_topology to PrimitiveTopology::LineList.
//! * Vertices and Indices describe the two points that make up a line.
use std::{borrow::Cow, iter, mem::size_of};
use std::{iter, mem::size_of};
use bytemuck::{Pod, Zeroable};
use wgpu::util::DeviceExt;
@ -156,10 +156,7 @@ impl crate::framework::Example for Example {
let sample_count = max_sample_count;
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: None,

5
examples/src/render_to_texture/mod.rs
View File

@ -28,10 +28,7 @@ async fn run(_path: Option<String>) {
.await
.unwrap();
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let render_target = device.create_texture(&wgpu::TextureDescriptor {
label: None,

7
examples/src/repeated_compute/mod.rs
View File

@ -180,12 +180,7 @@ impl WgpuContext {
.unwrap();
// Our shader, kindly compiled with Naga.
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!(
"shader.wgsl"
))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
// This is where the GPU will read from and write to.
let storage_buffer = device.create_buffer(&wgpu::BufferDescriptor {

7
examples/src/shadow/mod.rs
View File

@ -1,4 +1,4 @@
use std::{borrow::Cow, f32::consts, iter, mem::size_of, ops::Range, sync::Arc};
use std::{f32::consts, iter, mem::size_of, ops::Range, sync::Arc};
use bytemuck::{Pod, Zeroable};
use wgpu::util::{align_to, DeviceExt};
@ -447,10 +447,7 @@ impl crate::framework::Example for Example {
attributes: &vertex_attr,
};
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let shadow_pass = {
let uniform_size = size_of::<GlobalUniforms>() as wgpu::BufferAddress;

7
examples/src/skybox/mod.rs
View File

@ -1,5 +1,5 @@
use bytemuck::{Pod, Zeroable};
use std::{borrow::Cow, f32::consts, mem::size_of};
use std::{f32::consts, mem::size_of};
use wgpu::{util::DeviceExt, AstcBlock, AstcChannel};
const IMAGE_SIZE: u32 = 256;
@ -168,10 +168,7 @@ impl crate::framework::Example for Example {
});
// Create the render pipeline
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let camera = Camera {
screen_size: (config.width, config.height),

7
examples/src/srgb_blend/mod.rs
View File

@ -1,5 +1,5 @@
use bytemuck::{Pod, Zeroable};
use std::{borrow::Cow, mem};
use std::mem;
use wgpu::util::DeviceExt;
#[repr(C)]
@ -103,10 +103,7 @@ impl<const SRGB: bool> crate::framework::Example for Example<SRGB> {
label: None,
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let vertex_buffers = [wgpu::VertexBufferLayout {
array_stride: vertex_size as wgpu::BufferAddress,

6
examples/src/stencil_triangles/mod.rs
View File

@ -1,5 +1,4 @@
use bytemuck::{Pod, Zeroable};
use std::borrow::Cow;
use std::mem::size_of;
use wgpu::util::DeviceExt;
@ -53,10 +52,7 @@ impl crate::framework::Example for Example {
push_constant_ranges: &[],
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let vertex_buffers = [wgpu::VertexBufferLayout {
array_stride: vertex_size as wgpu::BufferAddress,

5
examples/src/storage_texture/mod.rs
View File

@ -44,10 +44,7 @@ async fn run(_path: Option<String>) {
.await
.unwrap();
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let storage_texture = device.create_texture(&wgpu::TextureDescriptor {
label: None,

7
examples/src/timestamp_queries/mod.rs
View File

@ -227,7 +227,7 @@ async fn run() {
let queries = submit_render_and_compute_pass_with_queries(&device, &queue);
let raw_results = queries.wait_for_results(&device);
println!("Raw timestamp buffer contents: {:?}", raw_results);
println!("Raw timestamp buffer contents: {raw_results:?}");
QueryResults::from_raw_results(raw_results, timestamps_inside_passes).print(&queue);
}
@ -239,10 +239,7 @@ fn submit_render_and_compute_pass_with_queries(
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
let mut queries = Queries::new(device, QueryResults::NUM_QUERIES);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!("shader.wgsl"))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
if device
.features()

7
examples/src/uniform_values/mod.rs
View File

@ -122,12 +122,7 @@ impl WgpuContext {
.await
.unwrap();
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!(
"shader.wgsl"
))),
});
let shader = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
// (2)
let uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor {

12
examples/src/water/mod.rs
View File

@ -3,7 +3,7 @@ mod point_gen;
use bytemuck::{Pod, Zeroable};
use glam::Vec3;
use nanorand::{Rng, WyRand};
use std::{borrow::Cow, f32::consts, iter, mem::size_of};
use std::{f32::consts, iter, mem::size_of};
use wgpu::util::DeviceExt;
///
@ -493,14 +493,8 @@ impl crate::framework::Example for Example {
});
// Upload/compile them to GPU code.
let terrain_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("terrain"),
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("terrain.wgsl"))),
});
let water_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("water"),
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("water.wgsl"))),
});
let terrain_module = device.create_shader_module(wgpu::include_wgsl!("terrain.wgsl"));
let water_module = device.create_shader_module(wgpu::include_wgsl!("water.wgsl"));
// Create the render pipelines. These describe how the data will flow through the GPU, and what
// constraints and modifiers it will have.

48
naga-cli/src/bin/naga.rs
View File

@ -465,6 +465,7 @@ fn run() -> anyhow::Result<()> {
let Parsed {
mut module,
input_text,
language,
} = parse_input(input_path, input, &params)?;
// Include debugging information if requested.
@ -477,6 +478,7 @@ fn run() -> anyhow::Result<()> {
params.spv_out.debug_info = Some(naga::back::spv::DebugInfo {
source_code: input_text,
file_name: input_path,
language,
})
} else {
eprintln!(
@ -579,6 +581,7 @@ fn run() -> anyhow::Result<()> {
struct Parsed {
module: naga::Module,
input_text: Option<String>,
language: naga::back::spv::SourceLanguage,
}
fn parse_input(input_path: &Path, input: Vec<u8>, params: &Parameters) -> anyhow::Result<Parsed> {
@ -593,16 +596,26 @@ fn parse_input(input_path: &Path, input: Vec<u8>, params: &Parameters) -> anyhow
.context("Unable to determine --input-kind from filename")?,
};
let (module, input_text) = match input_kind {
InputKind::Bincode => (bincode::deserialize(&input)?, None),
InputKind::SpirV => {
naga::front::spv::parse_u8_slice(&input, &params.spv_in).map(|m| (m, None))?
}
Ok(match input_kind {
InputKind::Bincode => Parsed {
module: bincode::deserialize(&input)?,
input_text: None,
language: naga::back::spv::SourceLanguage::Unknown,
},
InputKind::SpirV => Parsed {
module: naga::front::spv::parse_u8_slice(&input, &params.spv_in)?,
input_text: None,
language: naga::back::spv::SourceLanguage::Unknown,
},
InputKind::Wgsl => {
let input = String::from_utf8(input)?;
let result = naga::front::wgsl::parse_str(&input);
match result {
Ok(v) => (v, Some(input)),
Ok(v) => Parsed {
module: v,
input_text: Some(input),
language: naga::back::spv::SourceLanguage::WGSL,
},
Err(ref e) => {
let message = anyhow!(
"Could not parse WGSL:\n{}",
@ -631,8 +644,8 @@ fn parse_input(input_path: &Path, input: Vec<u8>, params: &Parameters) -> anyhow
};
let input = String::from_utf8(input)?;
let mut parser = naga::front::glsl::Frontend::default();
(
parser
Parsed {
module: parser
.parse(
&naga::front::glsl::Options {
stage: shader_stage.0,
@ -649,12 +662,11 @@ fn parse_input(input_path: &Path, input: Vec<u8>, params: &Parameters) -> anyhow
error.emit_to_writer_with_path(&mut writer, &input, filename);
std::process::exit(1);
}),
Some(input),
)
input_text: Some(input),
language: naga::back::spv::SourceLanguage::GLSL,
}
}
};
Ok(Parsed { module, input_text })
})
}
fn write_output(
@ -833,11 +845,15 @@ fn bulk_validate(args: Args, params: &Parameters) -> anyhow::Result<()> {
let path = Path::new(&input_path);
let input = fs::read(path)?;
let Parsed { module, input_text } = match parse_input(path, input, params) {
let Parsed {
module,
input_text,
language: _,
} = match parse_input(path, input, params) {
Ok(parsed) => parsed,
Err(error) => {
invalid.push(input_path.clone());
eprintln!("Error validating {}:", input_path);
eprintln!("Error validating {input_path}:");
eprintln!("{error}");
continue;
}
@ -850,7 +866,7 @@ fn bulk_validate(args: Args, params: &Parameters) -> anyhow::Result<()> {
if let Err(error) = validator.validate(&module) {
invalid.push(input_path.clone());
eprintln!("Error validating {}:", input_path);
eprintln!("Error validating {input_path}:");
if let Some(input) = &input_text {
let filename = path.file_name().and_then(std::ffi::OsStr::to_str);
emit_annotated_error(&error, filename.unwrap_or("input"), input);

2
naga/Cargo.toml
View File

@ -81,7 +81,7 @@ serde = { version = "1.0.210", features = ["derive"], optional = true }
petgraph = { version = "0.6", optional = true }
pp-rs = { version = "0.2.1", optional = true }
hexf-parse = { version = "0.2.1", optional = true }
unicode-xid = { version = "0.2.5", optional = true }
unicode-xid = { version = "0.2.6", optional = true }
[build-dependencies]
cfg_aliases.workspace = true

2
naga/src/back/dot/mod.rs
View File

@ -698,7 +698,7 @@ fn write_function_expressions(
E::RayQueryGetIntersection { query, committed } => {
edges.insert("", query);
let ty = if committed { "Committed" } else { "Candidate" };
(format!("rayQueryGet{}Intersection", ty).into(), 4)
(format!("rayQueryGet{ty}Intersection").into(), 4)
}
E::SubgroupBallotResult => ("SubgroupBallotResult".into(), 4),
E::SubgroupOperationResult { .. } => ("SubgroupOperationResult".into(), 4),

12
naga/src/back/glsl/mod.rs
View File

@ -2645,15 +2645,15 @@ impl<'a, W: Write> Writer<'a, W> {
match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero which is needed for a valid glsl float constant
crate::Literal::F64(value) => write!(self.out, "{:?}LF", value)?,
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
crate::Literal::F64(value) => write!(self.out, "{value:?}LF")?,
crate::Literal::F32(value) => write!(self.out, "{value:?}")?,
// Unsigned integers need a `u` at the end
//
// While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
// always write it as the extra branch wouldn't have any benefit in readability
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
crate::Literal::U32(value) => write!(self.out, "{value}u")?,
crate::Literal::I32(value) => write!(self.out, "{value}")?,
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
crate::Literal::I64(_) => {
return Err(Error::Custom("GLSL has no 64-bit integer type".into()));
}
@ -4614,7 +4614,7 @@ impl<'a, W: Write> Writer<'a, W> {
for i in 0..count.get() {
// Add the array accessor and recurse.
segments.push(format!("[{}]", i));
segments.push(format!("[{i}]"));
self.collect_push_constant_items(base, segments, layouter, offset, items);
segments.pop();
}

3
naga/src/back/hlsl/help.rs
View File

@ -1046,8 +1046,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
}
ref other => {
return Err(super::Error::Unimplemented(format!(
"Array length of base {:?}",
other
"Array length of base {other:?}"
)))
}
};

2
naga/src/back/hlsl/storage.rs
View File

@ -350,7 +350,7 @@ impl<W: fmt::Write> super::Writer<'_, W> {
self.write_store_value(module, &value, func_ctx)?;
writeln!(self.out, "));")?;
} else {
write!(self.out, "{}{}.Store(", level, var_name)?;
write!(self.out, "{level}{var_name}.Store(")?;
self.write_storage_address(module, &chain, func_ctx)?;
write!(self.out, ", ")?;
self.write_store_value(module, &value, func_ctx)?;

12
naga/src/back/hlsl/writer.rs
View File

@ -965,7 +965,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
let constant = &module.constants[handle];
self.write_type(module, constant.ty)?;
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, " {}", name)?;
write!(self.out, " {name}")?;
// Write size for array type
if let TypeInner::Array { base, size, .. } = module.types[constant.ty].inner {
self.write_array_size(module, base, size)?;
@ -2383,11 +2383,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
// decimal part even it's zero
crate::Literal::F64(value) => write!(self.out, "{value:?}L")?,
crate::Literal::F32(value) => write!(self.out, "{value:?}")?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::U64(value) => write!(self.out, "{}uL", value)?,
crate::Literal::I64(value) => write!(self.out, "{}L", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
crate::Literal::U32(value) => write!(self.out, "{value}u")?,
crate::Literal::I32(value) => write!(self.out, "{value}")?,
crate::Literal::U64(value) => write!(self.out, "{value}uL")?,
crate::Literal::I64(value) => write!(self.out, "{value}L")?,
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
return Err(Error::Custom(
"Abstract types should not appear in IR presented to backends".into(),

3
naga/src/back/msl/mod.rs
View File

@ -437,8 +437,7 @@ impl Options {
})
}
LocationMode::Uniform => Err(Error::GenericValidation(format!(
"Unexpected Binding::Location({}) for the Uniform mode",
location
"Unexpected Binding::Location({location}) for the Uniform mode"
))),
},
}

149
naga/src/back/msl/writer.rs
View File

@ -376,6 +376,11 @@ pub struct Writer<W> {
/// Set of (struct type, struct field index) denoting which fields require
/// padding inserted **before** them (i.e. between fields at index - 1 and index)
struct_member_pads: FastHashSet<(Handle<crate::Type>, u32)>,
/// Name of the loop reachability macro.
///
/// See `emit_loop_reachable_macro` for details.
loop_reachable_macro_name: String,
}
impl crate::Scalar {
@ -665,6 +670,7 @@ impl<W: Write> Writer<W> {
#[cfg(test)]
put_block_stack_pointers: Default::default(),
struct_member_pads: FastHashSet::default(),
loop_reachable_macro_name: String::default(),
}
}
@ -675,6 +681,128 @@ impl<W: Write> Writer<W> {
self.out
}
/// Define a macro to invoke before loops, to defeat MSL infinite loop
/// reasoning.
///
/// If we haven't done so already, emit the definition of a preprocessor
/// macro to be invoked before each loop in the generated MSL, to ensure
/// that the MSL compiler's optimizations do not remove bounds checks.
///
/// Only the first call to this function for a given module actually causes
/// the macro definition to be written. Subsequent loops can simply use the
/// prior macro definition, since macros aren't block-scoped.
///
/// # What is this trying to solve?
///
/// In Metal Shading Language, an infinite loop has undefined behavior.
/// (This rule is inherited from C++14.) This means that, if the MSL
/// compiler determines that a given loop will never exit, it may assume
/// that it is never reached. It may thus assume that any conditions
/// sufficient to cause the loop to be reached must be false. Like many
/// optimizing compilers, MSL uses this kind of analysis to establish limits
/// on the range of values variables involved in those conditions might
/// hold.
///
/// For example, suppose the MSL compiler sees the code:
///
/// ```ignore
/// if (i >= 10) {
/// while (true) { }
/// }
/// ```
///
/// It will recognize that the `while` loop will never terminate, conclude
/// that it must be unreachable, and thus infer that, if this code is
/// reached, then `i < 10` at that point.
///
/// Now suppose that, at some point where `i` has the same value as above,
/// the compiler sees the code:
///
/// ```ignore
/// if (i < 10) {
/// a[i] = 1;
/// }
/// ```
///
/// Because the compiler is confident that `i < 10`, it will make the
/// assignment to `a[i]` unconditional, rewriting this code as, simply:
///
/// ```ignore
/// a[i] = 1;
/// ```
///
/// If that `if` condition was injected by Naga to implement a bounds check,
/// the MSL compiler's optimizations could allow out-of-bounds array
/// accesses to occur.
///
/// Naga cannot feasibly anticipate whether the MSL compiler will determine
/// that a loop is infinite, so an attacker could craft a Naga module
/// containing an infinite loop protected by conditions that cause the Metal
/// compiler to remove bounds checks that Naga injected elsewhere in the
/// function.
///
/// This rewrite could occur even if the conditional assignment appears
/// *before* the `while` loop, as long as `i < 10` by the time the loop is
/// reached. This would allow the attacker to save the results of
/// unauthorized reads somewhere accessible before entering the infinite
/// loop. But even worse, the MSL compiler has been observed to simply
/// delete the infinite loop entirely, so that even code dominated by the
/// loop becomes reachable. This would make the attack even more flexible,
/// since shaders that would appear to never terminate would actually exit
/// nicely, after having stolen data from elsewhere in the GPU address
/// space.
///
/// Ideally, Naga would prevent UB entirely via some means that persuades
/// the MSL compiler that no loop Naga generates is infinite. One approach
/// would be to add inline assembly to each loop that is annotated as
/// potentially branching out of the loop, but which in fact generates no
/// instructions. Unfortunately, inline assembly is not handled correctly by
/// some Metal device drivers. Further experimentation hasn't produced a
/// satisfactory approach.
///
/// Instead, we accept that the MSL compiler may determine that some loops
/// are infinite, and focus instead on preventing the range analysis from
/// being affected. We transform *every* loop into something like this:
///
/// ```ignore
/// if (volatile bool unpredictable = true; unpredictable)
/// while (true) { }
/// ```
///
/// Since the `volatile` qualifier prevents the compiler from assuming that
/// the `if` condition is true, it cannot be sure the infinite loop is
/// reached, and thus it cannot assume the entire structure is unreachable.
/// This prevents the range analysis impact described above.
///
/// Unfortunately, what makes this a kludge, not a hack, is that this
/// solution leaves the GPU executing a pointless conditional branch, at
/// runtime, before each loop. There's no part of the system that has a
/// global enough view to be sure that `unpredictable` is true, and remove
/// it from the code.
///
/// To make our output a bit more legible, we pull the condition out into a
/// preprocessor macro defined at the top of the module.
///
/// This approach is also used by Chromium WebGPU's Dawn shader compiler, as of
/// <https://github.com/google/dawn/commit/ffd485c685040edb1e678165dcbf0e841cfa0298>.
fn emit_loop_reachable_macro(&mut self) -> BackendResult {
if !self.loop_reachable_macro_name.is_empty() {
return Ok(());
}
self.loop_reachable_macro_name = self.namer.call("LOOP_IS_REACHABLE");
let loop_reachable_volatile_name = self.namer.call("unpredictable_jump_over_loop");
writeln!(
self.out,
"#define {} if (volatile bool {} = true; {})",
self.loop_reachable_macro_name,
loop_reachable_volatile_name,
loop_reachable_volatile_name,
)?;
Ok(())
}
fn put_call_parameters(
&mut self,
parameters: impl Iterator<Item = Handle<crate::Expression>>,
@ -2924,10 +3052,15 @@ impl<W: Write> Writer<W> {
ref continuing,
break_if,
} => {
self.emit_loop_reachable_macro()?;
if !continuing.is_empty() || break_if.is_some() {
let gate_name = self.namer.call("loop_init");
writeln!(self.out, "{level}bool {gate_name} = true;")?;
writeln!(self.out, "{level}while(true) {{")?;
writeln!(
self.out,
"{level}{} while(true) {{",
self.loop_reachable_macro_name,
)?;
let lif = level.next();
let lcontinuing = lif.next();
writeln!(self.out, "{lif}if (!{gate_name}) {{")?;
@ -2942,7 +3075,11 @@ impl<W: Write> Writer<W> {
writeln!(self.out, "{lif}}}")?;
writeln!(self.out, "{lif}{gate_name} = false;")?;
} else {
writeln!(self.out, "{level}while(true) {{")?;
writeln!(
self.out,
"{level}{} while(true) {{",
self.loop_reachable_macro_name,
)?;
}
self.put_block(level.next(), body, context)?;
writeln!(self.out, "{level}}}")?;
@ -3379,6 +3516,7 @@ impl<W: Write> Writer<W> {
&[CLAMPED_LOD_LOAD_PREFIX],
&mut self.names,
);
self.loop_reachable_macro_name.clear();
self.struct_member_pads.clear();
writeln!(
@ -3682,12 +3820,11 @@ impl<W: Write> Writer<W> {
writeln!(self.out)?;
writeln!(
self.out,
"{} {defined_func_name}({arg_type_name} arg) {{
"{struct_name} {defined_func_name}({arg_type_name} arg) {{
{other_type_name} other;
{arg_type_name} fract = {NAMESPACE}::{called_func_name}(arg, other);
return {}{{ fract, other }};
}}",
struct_name, struct_name
return {struct_name}{{ fract, other }};
}}"
)?;
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}

5
naga/src/back/pipeline_constants.rs
View File

@ -14,7 +14,10 @@ use thiserror::Error;
pub enum PipelineConstantError {
#[error("Missing value for pipeline-overridable constant with identifier string: '{0}'")]
MissingValue(String),
#[error("Source f64 value needs to be finite (NaNs and Inifinites are not allowed) for number destinations")]
#[error(
"Source f64 value needs to be finite ({}) for number destinations",
"NaNs and Inifinites are not allowed"
)]
SrcNeedsToBeFinite,
#[error("Source f64 value doesn't fit in destination")]
DstRangeTooSmall,

159
naga/src/back/spv/block.rs
View File

@ -4,8 +4,7 @@ Implementations for `BlockContext` methods.
use super::{
helpers, index::BoundsCheckResult, make_local, selection::Selection, Block, BlockContext,
Dimension, Error, Instruction, LocalType, LookupType, LoopContext, ResultMember, Writer,
WriterFlags,
Dimension, Error, Instruction, LocalType, LookupType, ResultMember, Writer, WriterFlags,
};
use crate::{arena::Handle, proc::TypeResolution, Statement};
use spirv::Word;
@ -39,7 +38,7 @@ enum ExpressionPointer {
}
/// The termination statement to be added to the end of the block
pub enum BlockExit {
enum BlockExit {
/// Generates an OpReturn (void return)
Return,
/// Generates an OpBranch to the specified block
@ -60,6 +59,36 @@ pub enum BlockExit {
},
}
/// What code generation did with a provided [`BlockExit`] value.
///
/// A function that accepts a [`BlockExit`] argument should return a value of
/// this type, to indicate whether the code it generated ended up using the
/// provided exit, or ignored it and did a non-local exit of some other kind
/// (say, [`Break`] or [`Continue`]). Some callers must use this information to
/// decide whether to generate the target block at all.
///
/// [`Break`]: Statement::Break
/// [`Continue`]: Statement::Continue
#[must_use]
enum BlockExitDisposition {
/// The generated code used the provided `BlockExit` value. If it included a
/// block label, the caller should be sure to actually emit the block it
/// refers to.
Used,
/// The generated code did not use the provided `BlockExit` value. If it
/// included a block label, the caller should not bother to actually emit
/// the block it refers to, unless it knows the block is needed for
/// something else.
Discarded,
}
#[derive(Clone, Copy, Default)]
struct LoopContext {
continuing_id: Option<Word>,
break_id: Option<Word>,
}
#[derive(Debug)]
pub(crate) struct DebugInfoInner<'a> {
pub source_code: &'a str,
@ -200,10 +229,7 @@ impl<'w> BlockContext<'w> {
fn is_intermediate(&self, expr_handle: Handle<crate::Expression>) -> bool {
match self.ir_function.expressions[expr_handle] {
crate::Expression::GlobalVariable(handle) => {
match self.ir_module.global_variables[handle].space {
crate::AddressSpace::Handle => false,
_ => true,
}
self.ir_module.global_variables[handle].space != crate::AddressSpace::Handle
}
crate::Expression::LocalVariable(_) => true,
crate::Expression::FunctionArgument(index) => {
@ -346,6 +372,32 @@ impl<'w> BlockContext<'w> {
load_id
}
crate::TypeInner::Array {
base: ty_element, ..
} => {
let index_id = self.cached[index];
let base_id = self.cached[base];
let base_ty = match self.fun_info[base].ty {
TypeResolution::Handle(handle) => handle,
TypeResolution::Value(_) => {
return Err(Error::Validation(
"Array types should always be in the arena",
))
}
};
let (id, variable) = self.writer.promote_access_expression_to_variable(
&self.ir_module.types,
result_type_id,
base_id,
base_ty,
index_id,
ty_element,
block,
)?;
self.function.internal_variables.push(variable);
id
}
// wgpu#4337: Support `crate::TypeInner::Matrix`
ref other => {
log::error!(
"Unable to access base {:?} of type {:?}",
@ -353,7 +405,7 @@ impl<'w> BlockContext<'w> {
other
);
return Err(Error::Validation(
"only vectors may be dynamically indexed by value",
"only vectors and arrays may be dynamically indexed by value",
));
}
}
@ -2037,14 +2089,30 @@ impl<'w> BlockContext<'w> {
}
}
pub(super) fn write_block(
/// Generate one or more SPIR-V blocks for `naga_block`.
///
/// Use `label_id` as the label for the SPIR-V entry point block.
///
/// If control reaches the end of the SPIR-V block, terminate it according
/// to `exit`. This function's return value indicates whether it acted on
/// this parameter or not; see [`BlockExitDisposition`].
///
/// If the block contains [`Break`] or [`Continue`] statements,
/// `loop_context` supplies the labels of the SPIR-V blocks to jump to. If
/// either of these labels are `None`, then it should have been a Naga
/// validation error for the corresponding statement to occur in this
/// context.
///
/// [`Break`]: Statement::Break
/// [`Continue`]: Statement::Continue
fn write_block(
&mut self,
label_id: Word,
naga_block: &crate::Block,
exit: BlockExit,
loop_context: LoopContext,
debug_info: Option<&DebugInfoInner>,
) -> Result<(), Error> {
) -> Result<BlockExitDisposition, Error> {
let mut block = Block::new(label_id);
for (statement, span) in naga_block.span_iter() {
if let (Some(debug_info), false) = (
@ -2080,7 +2148,7 @@ impl<'w> BlockContext<'w> {
self.function.consume(block, Instruction::branch(scope_id));
let merge_id = self.gen_id();
self.write_block(
let merge_used = self.write_block(
scope_id,
block_statements,
BlockExit::Branch { target: merge_id },
@ -2088,7 +2156,14 @@ impl<'w> BlockContext<'w> {
debug_info,
)?;
block = Block::new(merge_id);
match merge_used {
BlockExitDisposition::Used => {
block = Block::new(merge_id);
}
BlockExitDisposition::Discarded => {
return Ok(BlockExitDisposition::Discarded);
}
}
}
Statement::If {
condition,
@ -2124,7 +2199,11 @@ impl<'w> BlockContext<'w> {
);
if let Some(block_id) = accept_id {
self.write_block(
// We can ignore the `BlockExitDisposition` returned here because,
// even if `merge_id` is not actually reachable, it is always
// referred to by the `OpSelectionMerge` instruction we emitted
// earlier.
let _ = self.write_block(
block_id,
accept,
BlockExit::Branch { target: merge_id },
@ -2133,7 +2212,11 @@ impl<'w> BlockContext<'w> {
)?;
}
if let Some(block_id) = reject_id {
self.write_block(
// We can ignore the `BlockExitDisposition` returned here because,
// even if `merge_id` is not actually reachable, it is always
// referred to by the `OpSelectionMerge` instruction we emitted
// earlier.
let _ = self.write_block(
block_id,
reject,
BlockExit::Branch { target: merge_id },
@ -2211,7 +2294,15 @@ impl<'w> BlockContext<'w> {
} else {
merge_id
};
self.write_block(
// We can ignore the `BlockExitDisposition` returned here because
// `case_finish_id` is always referred to by either:
//
// - the `OpSwitch`, if it's the next case's label for a
// fall-through, or
//
// - the `OpSelectionMerge`, if it's the switch's overall merge
// block because there's no fall-through.
let _ = self.write_block(
*label_id,
&case.body,
BlockExit::Branch {
@ -2257,7 +2348,10 @@ impl<'w> BlockContext<'w> {
));
self.function.consume(block, Instruction::branch(body_id));
self.write_block(
// We can ignore the `BlockExitDisposition` returned here because,
// even if `continuing_id` is not actually reachable, it is always
// referred to by the `OpLoopMerge` instruction we emitted earlier.
let _ = self.write_block(
body_id,
body,
BlockExit::Branch {
@ -2280,7 +2374,10 @@ impl<'w> BlockContext<'w> {
},
};
self.write_block(
// We can ignore the `BlockExitDisposition` returned here because,
// even if `merge_id` is not actually reachable, it is always referred
// to by the `OpLoopMerge` instruction we emitted earlier.
let _ = self.write_block(
continuing_id,
continuing,
exit,
@ -2296,14 +2393,14 @@ impl<'w> BlockContext<'w> {
Statement::Break => {
self.function
.consume(block, Instruction::branch(loop_context.break_id.unwrap()));
return Ok(());
return Ok(BlockExitDisposition::Discarded);
}
Statement::Continue => {
self.function.consume(
block,
Instruction::branch(loop_context.continuing_id.unwrap()),
);
return Ok(());
return Ok(BlockExitDisposition::Discarded);
}
Statement::Return { value: Some(value) } => {
let value_id = self.cached[value];
@ -2322,15 +2419,15 @@ impl<'w> BlockContext<'w> {
None => Instruction::return_value(value_id),
};
self.function.consume(block, instruction);
return Ok(());
return Ok(BlockExitDisposition::Discarded);
}
Statement::Return { value: None } => {
self.function.consume(block, Instruction::return_void());
return Ok(());
return Ok(BlockExitDisposition::Discarded);
}
Statement::Kill => {
self.function.consume(block, Instruction::kill());
return Ok(());
return Ok(BlockExitDisposition::Discarded);
}
Statement::Barrier(flags) => {
self.writer.write_barrier(flags, &mut block);
@ -2696,6 +2793,24 @@ impl<'w> BlockContext<'w> {
};
self.function.consume(block, termination);
Ok(BlockExitDisposition::Used)
}
pub(super) fn write_function_body(
&mut self,
entry_id: Word,
debug_info: Option<&DebugInfoInner>,
) -> Result<(), Error> {
// We can ignore the `BlockExitDisposition` returned here because
// `BlockExit::Return` doesn't refer to a block.
let _ = self.write_block(
entry_id,
&self.ir_function.body,
super::block::BlockExit::Return,
LoopContext::default(),
debug_info,
)?;
Ok(())
}
}

15
naga/src/back/spv/helpers.rs
View File

@ -85,20 +85,9 @@ impl crate::AddressSpace {
/// Return true if the global requires a type decorated with `Block`.
///
/// Vulkan spec v1.3 §15.6.2, "Descriptor Set Interface", says:
/// See [`back::spv::GlobalVariable`] for details.
///
/// > Variables identified with the `Uniform` storage class are used to
/// > access transparent buffer backed resources. Such variables must
/// > be:
/// >
/// > - typed as `OpTypeStruct`, or an array of this type,
/// >
/// > - identified with a `Block` or `BufferBlock` decoration, and
/// >
/// > - laid out explicitly using the `Offset`, `ArrayStride`, and
/// > `MatrixStride` decorations as specified in §15.6.4, "Offset
/// > and Stride Assignment."
// See `back::spv::GlobalVariable::access_id` for details.
/// [`back::spv::GlobalVariable`]: super::GlobalVariable
pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariable) -> bool {
match var.space {
crate::AddressSpace::Uniform

214
naga/src/back/spv/index.rs
View File

@ -11,16 +11,31 @@ use crate::{arena::Handle, proc::BoundsCheckPolicy};
/// The results of performing a bounds check.
///
/// On success, `write_bounds_check` returns a value of this type.
/// On success, [`write_bounds_check`](BlockContext::write_bounds_check)
/// returns a value of this type. The caller can assume that the right
/// policy has been applied, and simply do what the variant says.
pub(super) enum BoundsCheckResult {
/// The index is statically known and in bounds, with the given value.
KnownInBounds(u32),
/// The given instruction computes the index to be used.
///
/// When [`BoundsCheckPolicy::Restrict`] is in force, this is a
/// clamped version of the index the user supplied.
///
/// When [`BoundsCheckPolicy::Unchecked`] is in force, this is
/// simply the index the user supplied. This variant indicates
/// that we couldn't prove statically that the index was in
/// bounds; otherwise we would have returned [`KnownInBounds`].
///
/// [`KnownInBounds`]: BoundsCheckResult::KnownInBounds
Computed(Word),
/// The given instruction computes a boolean condition which is true
/// if the index is in bounds.
///
/// This is returned when [`BoundsCheckPolicy::ReadZeroSkipWrite`]
/// is in force.
Conditional(Word),
}
@ -38,98 +53,163 @@ impl<'w> BlockContext<'w> {
///
/// Given `array`, an expression referring a runtime-sized array, return the
/// instruction id for the array's length.
///
/// Runtime-sized arrays may only appear in the values of global
/// variables, which must have one of the following Naga types:
///
/// 1. A runtime-sized array.
/// 2. A struct whose last member is a runtime-sized array.
/// 3. A binding array of 2.
///
/// Thus, the expression `array` has the form of:
///
/// - An optional [`AccessIndex`], for case 2, applied to...
/// - An optional [`Access`] or [`AccessIndex`], for case 3, applied to...
/// - A [`GlobalVariable`].
///
/// The generated SPIR-V takes into account wrapped globals; see
/// [`back::spv::GlobalVariable`] for details.
///
/// [`GlobalVariable`]: crate::Expression::GlobalVariable
/// [`AccessIndex`]: crate::Expression::AccessIndex
/// [`Access`]: crate::Expression::Access
/// [`base`]: crate::Expression::Access::base
/// [`back::spv::GlobalVariable`]: super::GlobalVariable
pub(super) fn write_runtime_array_length(
&mut self,
array: Handle<crate::Expression>,
block: &mut Block,
) -> Result<Word, Error> {
// Naga IR permits runtime-sized arrays as global variables, or as the
// final member of a struct that is a global variable, or one of these
// inside a buffer that is itself an element in a buffer bindings array.
// SPIR-V requires that runtime-sized arrays are wrapped in structs.
// See `helpers::global_needs_wrapper` and its uses.
let (opt_array_index_id, global_handle, opt_last_member_index) = match self
.ir_function
.expressions[array]
{
// The index into the binding array, if any.
let binding_array_index_id: Option<Word>;
// The handle to the Naga IR global we're referring to.
let global_handle: Handle<crate::GlobalVariable>;
// At the Naga type level, if the runtime-sized array is the final member of a
// struct, this is that member's index.
//
// This does not cover wrappers: if this backend wrapped the Naga global's
// type in a synthetic SPIR-V struct (see `global_needs_wrapper`), this is
// `None`.
let opt_last_member_index: Option<u32>;
// Inspect `array` and decide whether we have a binding array and/or an
// enclosing struct.
match self.ir_function.expressions[array] {
crate::Expression::AccessIndex { base, index } => {
match self.ir_function.expressions[base] {
// The global variable is an array of buffer bindings of structs,
// we are accessing one of them with a static index,
// and the last member of it.
crate::Expression::AccessIndex {
base: base_outer,
index: index_outer,
} => match self.ir_function.expressions[base_outer] {
// An `AccessIndex` of an `AccessIndex` must be a
// binding array holding structs whose last members are
// runtime-sized arrays.
crate::Expression::GlobalVariable(handle) => {
let index_id = self.get_index_constant(index_outer);
(Some(index_id), handle, Some(index))
binding_array_index_id = Some(index_id);
global_handle = handle;
opt_last_member_index = Some(index);
}
_ => {
return Err(Error::Validation(
"array length expression: AccessIndex(AccessIndex(Global))",
))
}
_ => return Err(Error::Validation("array length expression case-1a")),
},
// The global variable is an array of buffer bindings of structs,
// we are accessing one of them with a dynamic index,
// and the last member of it.
crate::Expression::Access {
base: base_outer,
index: index_outer,
} => match self.ir_function.expressions[base_outer] {
// Similarly, an `AccessIndex` of an `Access` must be a
// binding array holding structs whose last members are
// runtime-sized arrays.
crate::Expression::GlobalVariable(handle) => {
let index_id = self.cached[index_outer];
(Some(index_id), handle, Some(index))
binding_array_index_id = Some(index_id);
global_handle = handle;
opt_last_member_index = Some(index);
}
_ => {
return Err(Error::Validation(
"array length expression: AccessIndex(Access(Global))",
))
}
_ => return Err(Error::Validation("array length expression case-1b")),
},
// The global variable is a buffer, and we are accessing the last member.
crate::Expression::GlobalVariable(handle) => {
let global = &self.ir_module.global_variables[handle];
match self.ir_module.types[global.ty].inner {
// The global variable is an array of buffer bindings of run-time arrays.
crate::TypeInner::BindingArray { .. } => (Some(index), handle, None),
// The global variable is a struct, and we are accessing the last member
_ => (None, handle, Some(index)),
}
// An outer `AccessIndex` applied directly to a
// `GlobalVariable`. Since binding arrays can only contain
// structs, this must be referring to the last member of a
// struct that is a runtime-sized array.
binding_array_index_id = None;
global_handle = handle;
opt_last_member_index = Some(index);
}
_ => {
return Err(Error::Validation(
"array length expression: AccessIndex(<unexpected>)",
))
}
_ => return Err(Error::Validation("array length expression case-1c")),
}
}
// The global variable is an array of buffer bindings of arrays.
crate::Expression::Access { base, index } => match self.ir_function.expressions[base] {
crate::Expression::GlobalVariable(handle) => {
let index_id = self.cached[index];
let global = &self.ir_module.global_variables[handle];
match self.ir_module.types[global.ty].inner {
crate::TypeInner::BindingArray { .. } => (Some(index_id), handle, None),
_ => return Err(Error::Validation("array length expression case-2a")),
}
}
_ => return Err(Error::Validation("array length expression case-2b")),
},
// The global variable is a run-time array.
crate::Expression::GlobalVariable(handle) => {
let global = &self.ir_module.global_variables[handle];
if !global_needs_wrapper(self.ir_module, global) {
return Err(Error::Validation("array length expression case-3"));
}
(None, handle, None)
// A direct reference to a global variable. This must hold the
// runtime-sized array directly.
binding_array_index_id = None;
global_handle = handle;
opt_last_member_index = None;
}
_ => return Err(Error::Validation("array length expression case-4")),
};
// The verifier should have checked this, but make sure the inspection above
// agrees with the type about whether a binding array is involved.
//
// Eventually we do want to support `binding_array<array<T>>`. This check
// ensures that whoever relaxes the validator will get an error message from
// us, not just bogus SPIR-V.
let global = &self.ir_module.global_variables[global_handle];
match (
&self.ir_module.types[global.ty].inner,
binding_array_index_id,
) {
(&crate::TypeInner::BindingArray { .. }, Some(_)) => {}
(_, None) => {}
_ => {
return Err(Error::Validation(
"array length expression: bad binding array inference",
))
}
}
// SPIR-V allows runtime-sized arrays to appear only as the last member of a
// struct. Determine this member's index.
let gvar = self.writer.global_variables[global_handle].clone();
let global = &self.ir_module.global_variables[global_handle];
let (last_member_index, gvar_id) = match opt_last_member_index {
Some(index) => (index, gvar.access_id),
None => {
if !global_needs_wrapper(self.ir_module, global) {
return Err(Error::Validation(
"pointer to a global that is not a wrapped array",
));
}
let needs_wrapper = global_needs_wrapper(self.ir_module, global);
let (last_member_index, gvar_id) = match (opt_last_member_index, needs_wrapper) {
(Some(index), false) => {
// At the Naga type level, the runtime-sized array appears as the
// final member of a struct, whose index is `index`. We didn't need to
// wrap this, since the Naga type meets SPIR-V's requirements already.
(index, gvar.access_id)
}
(None, true) => {
// At the Naga type level, the runtime-sized array does not appear
// within a struct. We wrapped this in an OpTypeStruct with nothing
// else in it, so the index is zero. OpArrayLength wants the pointer
// to the wrapper struct, so use `gvar.var_id`.
(0, gvar.var_id)
}
_ => {
return Err(Error::Validation(
"array length expression: bad SPIR-V wrapper struct inference",
));
}
};
let structure_id = match opt_array_index_id {
let structure_id = match binding_array_index_id {
// We are indexing inside a binding array, generate the access op.
Some(index_id) => {
let element_type_id = match self.ir_module.types[global.ty].inner {
@ -293,6 +373,8 @@ impl<'w> BlockContext<'w> {
/// Write an index bounds comparison to `block`, if needed.
///
/// This is used to implement [`BoundsCheckPolicy::ReadZeroSkipWrite`].
///
/// If we're able to determine statically that `index` is in bounds for
/// `sequence`, return `KnownInBounds(value)`, where `value` is the actual
/// value of the index. (In principle, one could know that the index is in
@ -413,11 +495,23 @@ impl<'w> BlockContext<'w> {
/// Emit code for bounds checks for an array, vector, or matrix access.
///
/// This implements either `index_bounds_check_policy` or
/// `buffer_bounds_check_policy`, depending on the address space of the
/// pointer being accessed.
/// This tries to handle all the critical steps for bounds checks:
///
/// Return a `BoundsCheckResult` indicating how the index should be
/// - First, select the appropriate bounds check policy for `base`,
/// depending on its address space.
///
/// - Next, analyze `index` to see if its value is known at
/// compile time, in which case we can decide statically whether
/// the index is in bounds.
///
/// - If the index's value is not known at compile time, emit code to:
///
/// - restrict its value (for [`BoundsCheckPolicy::Restrict`]), or
///
/// - check whether it's in bounds (for
/// [`BoundsCheckPolicy::ReadZeroSkipWrite`]).
///
/// Return a [`BoundsCheckResult`] indicating how the index should be
/// consumed. See that type's documentation for details.
pub(super) fn write_bounds_check(
&mut self,

87
naga/src/back/spv/mod.rs
View File

@ -16,7 +16,7 @@ mod selection;
mod subgroup;
mod writer;
pub use spirv::Capability;
pub use spirv::{Capability, SourceLanguage};
use crate::arena::{Handle, HandleVec};
use crate::proc::{BoundsCheckPolicies, TypeResolution};
@ -89,6 +89,7 @@ impl IdGenerator {
pub struct DebugInfo<'a> {
pub source_code: &'a str,
pub file_name: &'a std::path::Path,
pub language: SourceLanguage,
}
/// A SPIR-V block to which we are still adding instructions.
@ -143,6 +144,7 @@ struct Function {
signature: Option<Instruction>,
parameters: Vec<FunctionArgument>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
internal_variables: Vec<LocalVariable>,
blocks: Vec<TerminatedBlock>,
entry_point_context: Option<EntryPointContext>,
}
@ -465,38 +467,75 @@ enum CachedConstant {
ZeroValue(Word),
}
/// The SPIR-V representation of a [`crate::GlobalVariable`].
///
/// In the Vulkan spec 1.3.296, the section [Descriptor Set Interface][dsi] says:
///
/// > Variables identified with the `Uniform` storage class are used to access
/// > transparent buffer backed resources. Such variables *must* be:
/// >
/// > - typed as `OpTypeStruct`, or an array of this type,
/// >
/// > - identified with a `Block` or `BufferBlock` decoration, and
/// >
/// > - laid out explicitly using the `Offset`, `ArrayStride`, and `MatrixStride`
/// > decorations as specified in "Offset and Stride Assignment".
///
/// This is followed by identical language for the `StorageBuffer`,
/// except that a `BufferBlock` decoration is not allowed.
///
/// When we encounter a global variable in the [`Storage`] or [`Uniform`]
/// address spaces whose type is not already [`Struct`], this backend implicitly
/// wraps the global variable in a struct: we generate a SPIR-V global variable
/// holding an `OpTypeStruct` with a single member, whose type is what the Naga
/// global's type would suggest, decorated as required above.
///
/// The [`helpers::global_needs_wrapper`] function determines whether a given
/// [`crate::GlobalVariable`] needs to be wrapped.
///
/// [dsi]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#interfaces-resources-descset
/// [`Storage`]: crate::AddressSpace::Storage
/// [`Uniform`]: crate::AddressSpace::Uniform
/// [`Struct`]: crate::TypeInner::Struct
#[derive(Clone)]
struct GlobalVariable {
/// ID of the OpVariable that declares the global.
/// The SPIR-V id of the `OpVariable` that declares the global.
///
/// If you need the variable's value, use [`access_id`] instead of this
/// field. If we wrapped the Naga IR `GlobalVariable`'s type in a struct to
/// comply with Vulkan's requirements, then this points to the `OpVariable`
/// with the synthesized struct type, whereas `access_id` points to the
/// field of said struct that holds the variable's actual value.
/// If this global has been implicitly wrapped in an `OpTypeStruct`, this id
/// refers to the wrapper, not the original Naga value it contains. If you
/// need the Naga value, use [`access_id`] instead of this field.
///
/// If this global is not implicitly wrapped, this is the same as
/// [`access_id`].
///
/// This is used to compute the `access_id` pointer in function prologues,
/// and used for `ArrayLength` expressions, which do need the struct.
/// and used for `ArrayLength` expressions, which need to pass the wrapper
/// struct.
///
/// [`access_id`]: GlobalVariable::access_id
var_id: Word,
/// For `AddressSpace::Handle` variables, this ID is recorded in the function
/// prelude block (and reset before every function) as `OpLoad` of the variable.
/// It is then used for all the global ops, such as `OpImageSample`.
/// The loaded value of a `AddressSpace::Handle` global variable.
///
/// If the current function uses this global variable, this is the id of an
/// `OpLoad` instruction in the function's prologue that loads its value.
/// (This value is assigned as we write the prologue code of each function.)
/// It is then used for all operations on the global, such as `OpImageSample`.
handle_id: Word,
/// Actual ID used to access this variable.
/// For wrapped buffer variables, this ID is `OpAccessChain` into the
/// wrapper. Otherwise, the same as `var_id`.
/// The SPIR-V id of a pointer to this variable's Naga IR value.
///
/// Vulkan requires that globals in the `StorageBuffer` and `Uniform` storage
/// classes must be structs with the `Block` decoration, but WGSL and Naga IR
/// make no such requirement. So for such variables, we generate a wrapper struct
/// type with a single element of the type given by Naga, generate an
/// `OpAccessChain` for that member in the function prelude, and use that pointer
/// to refer to the global in the function body. This is the id of that access,
/// updated for each function in `write_function`.
/// If the current function uses this global variable, and it has been
/// implicitly wrapped in an `OpTypeStruct`, this is the id of an
/// `OpAccessChain` instruction in the function's prologue that refers to
/// the wrapped value inside the struct. (This value is assigned as we write
/// the prologue code of each function.) If you need the wrapper struct
/// itself, use [`var_id`] instead of this field.
///
/// If this global is not implicitly wrapped, this is the same as
/// [`var_id`].
///
/// [`var_id`]: GlobalVariable::var_id
access_id: Word,
}
@ -626,12 +665,6 @@ impl BlockContext<'_> {
}
}
#[derive(Clone, Copy, Default)]
struct LoopContext {
continuing_id: Option<Word>,
break_id: Option<Word>,
}
pub struct Writer {
physical_layout: PhysicalLayout,
logical_layout: LogicalLayout,

65
naga/src/back/spv/writer.rs
View File

@ -3,7 +3,7 @@ use super::{
helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
make_local, Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo,
EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction,
LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options,
LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, Options,
PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
};
use crate::{
@ -32,6 +32,9 @@ impl Function {
for local_var in self.variables.values() {
local_var.instruction.to_words(sink);
}
for internal_var in self.internal_variables.iter() {
internal_var.instruction.to_words(sink);
}
}
for instruction in block.body.iter() {
instruction.to_words(sink);
@ -135,6 +138,56 @@ impl Writer {
self.capabilities_used.insert(spirv::Capability::Shader);
}
#[allow(clippy::too_many_arguments)]
pub(super) fn promote_access_expression_to_variable(
&mut self,
ir_types: &UniqueArena<crate::Type>,
result_type_id: Word,
container_id: Word,
container_ty: Handle<crate::Type>,
index_id: Word,
element_ty: Handle<crate::Type>,
block: &mut Block,
) -> Result<(Word, LocalVariable), Error> {
let pointer_type_id =
self.get_pointer_id(ir_types, container_ty, spirv::StorageClass::Function)?;
let variable = {
let id = self.id_gen.next();
LocalVariable {
id,
instruction: Instruction::variable(
pointer_type_id,
id,
spirv::StorageClass::Function,
None,
),
}
};
block
.body
.push(Instruction::store(variable.id, container_id, None));
let element_pointer_id = self.id_gen.next();
let element_pointer_type_id =
self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?;
block.body.push(Instruction::access_chain(
element_pointer_type_id,
element_pointer_id,
variable.id,
&[index_id],
));
let id = self.id_gen.next();
block.body.push(Instruction::load(
result_type_id,
id,
element_pointer_id,
None,
));
Ok((id, variable))
}
/// Indicate that the code requires any one of the listed capabilities.
///
/// If nothing in `capabilities` appears in the available capabilities
@ -703,13 +756,7 @@ impl Writer {
next_id
};
context.write_block(
main_id,
&ir_function.body,
super::block::BlockExit::Return,
LoopContext::default(),
debug_info.as_ref(),
)?;
context.write_function_body(main_id, debug_info.as_ref())?;
// Consume the `BlockContext`, ending its borrows and letting the
// `Writer` steal back its cached expression table and temp_list.
@ -1967,7 +2014,7 @@ impl Writer {
source_file_id,
});
self.debugs.append(&mut Instruction::source_auto_continued(
spirv::SourceLanguage::Unknown,
debug_info.language,
0,
&debug_info_inner,
));

18
naga/src/back/wgsl/writer.rs
View File

@ -1221,31 +1221,31 @@ impl<W: Write> Writer<W> {
match expressions[expr] {
Expression::Literal(literal) => match literal {
crate::Literal::F32(value) => write!(self.out, "{}f", value)?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::F32(value) => write!(self.out, "{value}f")?,
crate::Literal::U32(value) => write!(self.out, "{value}u")?,
crate::Literal::I32(value) => {
// `-2147483648i` is not valid WGSL. The most negative `i32`
// value can only be expressed in WGSL using AbstractInt and
// a unary negation operator.
if value == i32::MIN {
write!(self.out, "i32({})", value)?;
write!(self.out, "i32({value})")?;
} else {
write!(self.out, "{}i", value)?;
write!(self.out, "{value}i")?;
}
}
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
crate::Literal::F64(value) => write!(self.out, "{:?}lf", value)?,
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
crate::Literal::I64(value) => {
// `-9223372036854775808li` is not valid WGSL. The most negative `i64`
// value can only be expressed in WGSL using AbstractInt and
// a unary negation operator.
if value == i64::MIN {
write!(self.out, "i64({})", value)?;
write!(self.out, "i64({value})")?;
} else {
write!(self.out, "{}li", value)?;
write!(self.out, "{value}li")?;
}
}
crate::Literal::U64(value) => write!(self.out, "{:?}lu", value)?,
crate::Literal::U64(value) => write!(self.out, "{value:?}lu")?,
crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
return Err(Error::Custom(
"Abstract types should not appear in IR presented to backends".into(),

8
naga/src/front/atomic_upgrade.rs
View File

@ -44,12 +44,8 @@ pub enum Error {
MultiMemberStruct,
#[error("encountered unsupported global initializer in an atomic variable")]
GlobalInitUnsupported,
}
impl From<Error> for crate::front::spv::Error {
fn from(source: Error) -> Self {
crate::front::spv::Error::AtomicUpgradeError(source)
}
#[error("expected to find a global variable")]
GlobalVariableMissing,
}
#[derive(Clone, Default)]

9
naga/src/front/glsl/context.rs
View File

@ -630,7 +630,8 @@ impl<'a> Context<'a> {
frontend.errors.push(Error {
kind: ErrorKind::SemanticError(
format!(
"Cannot apply operation to {left_inner:?} and {right_inner:?}"
"Cannot apply operation to {:?} and {:?}",
left_inner, right_inner
)
.into(),
),
@ -828,7 +829,8 @@ impl<'a> Context<'a> {
frontend.errors.push(Error {
kind: ErrorKind::SemanticError(
format!(
"Cannot apply operation to {left_inner:?} and {right_inner:?}"
"Cannot apply operation to {:?} and {:?}",
left_inner, right_inner
)
.into(),
),
@ -908,7 +910,8 @@ impl<'a> Context<'a> {
frontend.errors.push(Error {
kind: ErrorKind::SemanticError(
format!(
"Cannot apply operation to {left_inner:?} and {right_inner:?}"
"Cannot apply operation to {:?} and {:?}",
left_inner, right_inner
)
.into(),
),

3
naga/src/front/glsl/functions.rs
View File

@ -634,7 +634,8 @@ impl Frontend {
self.errors.push(Error {
kind: ErrorKind::SemanticError(
format!(
"'{name}': image needs {overload_access:?} access but only {call_access:?} was provided"
"'{}': image needs {:?} access but only {:?} was provided",
name, overload_access, call_access
)
.into(),
),

8
naga/src/front/glsl/parser/expressions.rs
View File

@ -38,7 +38,13 @@ impl<'source> ParsingContext<'source> {
TokenValue::FloatConstant(float) => {
if float.width != 32 {
frontend.errors.push(Error {
kind: ErrorKind::SemanticError("Unsupported floating-point value (expected single-precision floating-point number)".into()),
kind: ErrorKind::SemanticError(
concat!(
"Unsupported floating-point value ",
"(expected single-precision floating-point number)"
)
.into(),
),
meta: token.meta,
});
}

19
naga/src/front/glsl/variables.rs
View File

@ -294,14 +294,17 @@ impl Frontend {
.any(|i| components[i..].contains(&components[i - 1]));
if not_unique {
self.errors.push(Error {
kind:
ErrorKind::SemanticError(
format!(
"swizzle cannot have duplicate components in left-hand-side expression for \"{name:?}\""
)
.into(),
),
meta ,
kind: ErrorKind::SemanticError(
format!(
concat!(
"swizzle cannot have duplicate components in ",
"left-hand-side expression for \"{:?}\""
),
name
)
.into(),
),
meta,
})
}
}

14
naga/src/front/spv/error.rs
View File

@ -47,7 +47,13 @@ pub enum Error {
UnsupportedBinaryOperator(spirv::Word),
#[error("Naga supports OpTypeRuntimeArray in the StorageBuffer storage class only")]
UnsupportedRuntimeArrayStorageClass,
#[error("unsupported matrix stride {stride} for a {columns}x{rows} matrix with scalar width={width}")]
#[error(
"unsupported matrix stride {} for a {}x{} matrix with scalar width={}",
stride,
columns,
rows,
width
)]
UnsupportedMatrixStride {
stride: u32,
columns: u8,
@ -159,3 +165,9 @@ impl Error {
String::from_utf8(writer.into_inner()).unwrap()
}
}
impl From<atomic_upgrade::Error> for Error {
fn from(source: atomic_upgrade::Error) -> Self {
Error::AtomicUpgradeError(source)
}
}

329
naga/src/front/spv/mod.rs
View File

@ -565,11 +565,15 @@ impl<'a> BlockContext<'a> {
/// Descend into the expression with the given handle, locating a contained
/// global variable.
///
/// If the expression doesn't actually refer to something in a global
/// variable, we can't upgrade its type in a way that Naga validation would
/// pass, so reject the input instead.
///
/// This is used to track atomic upgrades.
fn get_contained_global_variable(
&self,
mut handle: Handle<crate::Expression>,
) -> Option<Handle<crate::GlobalVariable>> {
) -> Result<Handle<crate::GlobalVariable>, Error> {
log::debug!("\t\tlocating global variable in {handle:?}");
loop {
match self.expressions[handle] {
@ -583,14 +587,16 @@ impl<'a> BlockContext<'a> {
}
crate::Expression::GlobalVariable(h) => {
log::debug!("\t\t found {h:?}");
return Some(h);
return Ok(h);
}
_ => {
break;
}
}
}
None
Err(Error::AtomicUpgradeError(
crate::front::atomic_upgrade::Error::GlobalVariableMissing,
))
}
}
@ -1323,6 +1329,109 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
))
}
/// Return the Naga [`Expression`] for `pointer_id`, and its referent [`Type`].
///
/// Return a [`Handle`] for a Naga [`Expression`] that holds the value of
/// the SPIR-V instruction `pointer_id`, along with the [`Type`] to which it
/// is a pointer.
///
/// This may entail spilling `pointer_id`'s value to a temporary:
/// see [`get_expr_handle`]'s documentation.
///
/// [`Expression`]: crate::Expression
/// [`Type`]: crate::Type
/// [`Handle`]: crate::Handle
/// [`get_expr_handle`]: Frontend::get_expr_handle
fn get_exp_and_base_ty_handles(
&self,
pointer_id: spirv::Word,
ctx: &mut BlockContext,
emitter: &mut crate::proc::Emitter,
block: &mut crate::Block,
body_idx: usize,
) -> Result<(Handle<crate::Expression>, Handle<crate::Type>), Error> {
log::trace!("\t\t\tlooking up pointer expr {:?}", pointer_id);
let p_lexp_handle;
let p_lexp_ty_id;
{
let lexp = self.lookup_expression.lookup(pointer_id)?;
p_lexp_handle = self.get_expr_handle(pointer_id, lexp, ctx, emitter, block, body_idx);
p_lexp_ty_id = lexp.type_id;
};
log::trace!("\t\t\tlooking up pointer type {pointer_id:?}");
let p_ty = self.lookup_type.lookup(p_lexp_ty_id)?;
let p_ty_base_id = p_ty.base_id.ok_or(Error::InvalidAccessType(p_lexp_ty_id))?;
log::trace!("\t\t\tlooking up pointer base type {p_ty_base_id:?} of {p_ty:?}");
let p_base_ty = self.lookup_type.lookup(p_ty_base_id)?;
Ok((p_lexp_handle, p_base_ty.handle))
}
#[allow(clippy::too_many_arguments)]
fn parse_atomic_expr_with_value(
&mut self,
inst: Instruction,
emitter: &mut crate::proc::Emitter,
ctx: &mut BlockContext,
block: &mut crate::Block,
block_id: spirv::Word,
body_idx: usize,
atomic_function: crate::AtomicFunction,
) -> Result<(), Error> {
inst.expect(7)?;
let start = self.data_offset;
let result_type_id = self.next()?;
let result_id = self.next()?;
let pointer_id = self.next()?;
let _scope_id = self.next()?;
let _memory_semantics_id = self.next()?;
let value_id = self.next()?;
let span = self.span_from_with_op(start);
let (p_lexp_handle, p_base_ty_handle) =
self.get_exp_and_base_ty_handles(pointer_id, ctx, emitter, block, body_idx)?;
log::trace!("\t\t\tlooking up value expr {value_id:?}");
let v_lexp_handle = self.lookup_expression.lookup(value_id)?.handle;
block.extend(emitter.finish(ctx.expressions));
// Create an expression for our result
let r_lexp_handle = {
let expr = crate::Expression::AtomicResult {
ty: p_base_ty_handle,
comparison: false,
};
let handle = ctx.expressions.append(expr, span);
self.lookup_expression.insert(
result_id,
LookupExpression {
handle,
type_id: result_type_id,
block_id,
},
);
handle
};
emitter.start(ctx.expressions);
// Create a statement for the op itself
let stmt = crate::Statement::Atomic {
pointer: p_lexp_handle,
fun: atomic_function,
value: v_lexp_handle,
result: Some(r_lexp_handle),
};
block.push(stmt, span);
// Store any associated global variables so we can upgrade their types later
self.upgrade_atomics
.insert(ctx.get_contained_global_variable(p_lexp_handle)?);
Ok(())
}
/// Add the next SPIR-V block's contents to `block_ctx`.
///
/// Except for the function's entry block, `block_id` should be the label of
@ -3985,35 +4094,91 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
);
emitter.start(ctx.expressions);
}
Op::AtomicIIncrement => {
Op::AtomicLoad => {
inst.expect(6)?;
let start = self.data_offset;
let span = self.span_from_with_op(start);
let result_type_id = self.next()?;
let result_id = self.next()?;
let pointer_id = self.next()?;
let _scope_id = self.next()?;
let _memory_semantics_id = self.next()?;
let span = self.span_from_with_op(start);
log::trace!("\t\t\tlooking up expr {:?}", pointer_id);
let (p_lexp_handle, p_lexp_ty_id) = {
let lexp = self.lookup_expression.lookup(pointer_id)?;
let handle = get_expr_handle!(pointer_id, &lexp);
(handle, lexp.type_id)
let p_lexp_handle =
get_expr_handle!(pointer_id, self.lookup_expression.lookup(pointer_id)?);
// Create an expression for our result
let expr = crate::Expression::Load {
pointer: p_lexp_handle,
};
let handle = ctx.expressions.append(expr, span);
self.lookup_expression.insert(
result_id,
LookupExpression {
handle,
type_id: result_type_id,
block_id,
},
);
log::trace!("\t\t\tlooking up type {pointer_id:?}");
let p_ty = self.lookup_type.lookup(p_lexp_ty_id)?;
let p_ty_base_id =
p_ty.base_id.ok_or(Error::InvalidAccessType(p_lexp_ty_id))?;
// Store any associated global variables so we can upgrade their types later
self.upgrade_atomics
.insert(ctx.get_contained_global_variable(p_lexp_handle)?);
}
Op::AtomicStore => {
inst.expect(5)?;
let start = self.data_offset;
let pointer_id = self.next()?;
let _scope_id = self.next()?;
let _memory_semantics_id = self.next()?;
let value_id = self.next()?;
let span = self.span_from_with_op(start);
log::trace!("\t\t\tlooking up base type {p_ty_base_id:?} of {p_ty:?}");
let p_base_ty = self.lookup_type.lookup(p_ty_base_id)?;
log::trace!("\t\t\tlooking up pointer expr {:?}", pointer_id);
let p_lexp_handle =
get_expr_handle!(pointer_id, self.lookup_expression.lookup(pointer_id)?);
log::trace!("\t\t\tlooking up value expr {:?}", pointer_id);
let v_lexp_handle =
get_expr_handle!(value_id, self.lookup_expression.lookup(value_id)?);
block.extend(emitter.finish(ctx.expressions));
// Create a statement for the op itself
let stmt = crate::Statement::Store {
pointer: p_lexp_handle,
value: v_lexp_handle,
};
block.push(stmt, span);
emitter.start(ctx.expressions);
// Store any associated global variables so we can upgrade their types later
self.upgrade_atomics
.insert(ctx.get_contained_global_variable(p_lexp_handle)?);
}
Op::AtomicIIncrement | Op::AtomicIDecrement => {
inst.expect(6)?;
let start = self.data_offset;
let result_type_id = self.next()?;
let result_id = self.next()?;
let pointer_id = self.next()?;
let _scope_id = self.next()?;
let _memory_semantics_id = self.next()?;
let span = self.span_from_with_op(start);
let (p_exp_h, p_base_ty_h) = self.get_exp_and_base_ty_handles(
pointer_id,
ctx,
&mut emitter,
&mut block,
body_idx,
)?;
block.extend(emitter.finish(ctx.expressions));
// Create an expression for our result
let r_lexp_handle = {
let expr = crate::Expression::AtomicResult {
ty: p_base_ty.handle,
ty: p_base_ty_h,
comparison: false,
};
let handle = ctx.expressions.append(expr, span);
@ -4027,22 +4192,26 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
);
handle
};
emitter.start(ctx.expressions);
// Create a literal "1" since WGSL lacks an increment operation
// Create a literal "1" to use as our value
let one_lexp_handle = make_index_literal(
ctx,
1,
&mut block,
&mut emitter,
p_base_ty.handle,
p_lexp_ty_id,
p_base_ty_h,
result_type_id,
span,
)?;
// Create a statement for the op itself
let stmt = crate::Statement::Atomic {
pointer: p_lexp_handle,
fun: crate::AtomicFunction::Add,
pointer: p_exp_h,
fun: match inst.op {
Op::AtomicIIncrement => crate::AtomicFunction::Add,
_ => crate::AtomicFunction::Subtract,
},
value: one_lexp_handle,
result: Some(r_lexp_handle),
};
@ -4050,8 +4219,38 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
// Store any associated global variables so we can upgrade their types later
self.upgrade_atomics
.extend(ctx.get_contained_global_variable(p_lexp_handle));
.insert(ctx.get_contained_global_variable(p_exp_h)?);
}
Op::AtomicExchange
| Op::AtomicIAdd
| Op::AtomicISub
| Op::AtomicSMin
| Op::AtomicUMin
| Op::AtomicSMax
| Op::AtomicUMax
| Op::AtomicAnd
| Op::AtomicOr
| Op::AtomicXor => self.parse_atomic_expr_with_value(
inst,
&mut emitter,
ctx,
&mut block,
block_id,
body_idx,
match inst.op {
Op::AtomicExchange => crate::AtomicFunction::Exchange { compare: None },
Op::AtomicIAdd => crate::AtomicFunction::Add,
Op::AtomicISub => crate::AtomicFunction::Subtract,
Op::AtomicSMin => crate::AtomicFunction::Min,
Op::AtomicUMin => crate::AtomicFunction::Min,
Op::AtomicSMax => crate::AtomicFunction::Max,
Op::AtomicUMax => crate::AtomicFunction::Max,
Op::AtomicAnd => crate::AtomicFunction::And,
Op::AtomicOr => crate::AtomicFunction::InclusiveOr,
_ => crate::AtomicFunction::ExclusiveOr,
},
)?,
_ => {
return Err(Error::UnsupportedInstruction(self.state, inst.op));
}
@ -5709,33 +5908,48 @@ mod test {
];
let _ = super::parse_u8_slice(&bin, &Default::default()).unwrap();
}
}
#[cfg(all(feature = "wgsl-in", wgsl_out))]
#[test]
fn atomic_i_inc() {
#[cfg(all(test, feature = "wgsl-in", wgsl_out))]
mod test_atomic {
fn atomic_test(bytes: &[u8]) {
let _ = env_logger::builder().is_test(true).try_init();
let bytes = include_bytes!("../../../tests/in/spv/atomic_i_increment.spv");
let m = super::parse_u8_slice(bytes, &Default::default()).unwrap();
let mut validator = crate::valid::Validator::new(
let m = crate::front::spv::parse_u8_slice(bytes, &Default::default()).unwrap();
let mut wgsl = String::new();
let mut should_panic = false;
for vflags in [
crate::valid::ValidationFlags::all(),
crate::valid::ValidationFlags::empty(),
Default::default(),
);
let info = match validator.validate(&m) {
Err(e) => {
log::error!("{}", e.emit_to_string(""));
return;
}
Ok(i) => i,
};
let wgsl =
crate::back::wgsl::write_string(&m, &info, crate::back::wgsl::WriterFlags::empty())
.unwrap();
log::info!("atomic_i_increment:\n{wgsl}");
] {
let mut validator = crate::valid::Validator::new(vflags, Default::default());
match validator.validate(&m) {
Err(e) => {
log::error!("SPIR-V validation {}", e.emit_to_string(""));
should_panic = true;
}
Ok(i) => {
wgsl = crate::back::wgsl::write_string(
&m,
&i,
crate::back::wgsl::WriterFlags::empty(),
)
.unwrap();
log::info!("wgsl-out:\n{wgsl}");
break;
}
};
}
if should_panic {
panic!("validation error");
}
let m = match crate::front::wgsl::parse_str(&wgsl) {
Ok(m) => m,
Err(e) => {
log::error!("{}", e.emit_to_string(&wgsl));
log::error!("round trip WGSL validation {}", e.emit_to_string(&wgsl));
panic!("invalid module");
}
};
@ -5746,4 +5960,35 @@ mod test {
panic!("invalid generated wgsl");
}
}
#[test]
fn atomic_i_inc() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_i_increment.spv"
));
}
#[test]
fn atomic_load_and_store() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_load_and_store.spv"
));
}
#[test]
fn atomic_exchange() {
atomic_test(include_bytes!("../../../tests/in/spv/atomic_exchange.spv"));
}
#[test]
fn atomic_i_decrement() {
atomic_test(include_bytes!(
"../../../tests/in/spv/atomic_i_decrement.spv"
));
}
#[test]
fn atomic_i_add_and_sub() {
atomic_test(include_bytes!("../../../tests/in/spv/atomic_i_add_sub.spv"));
}
}

188
naga/src/front/wgsl/error.rs
View File

@ -298,32 +298,42 @@ impl<'a> Error<'a> {
match *self {
Error::Unexpected(unexpected_span, expected) => {
let expected_str = match expected {
ExpectedToken::Token(token) => {
match token {
Token::Separator(c) => format!("'{c}'"),
Token::Paren(c) => format!("'{c}'"),
Token::Attribute => "@".to_string(),
Token::Number(_) => "number".to_string(),
Token::Word(s) => s.to_string(),
Token::Operation(c) => format!("operation ('{c}')"),
Token::LogicalOperation(c) => format!("logical operation ('{c}')"),
Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"),
Token::AssignmentOperation(c) if c=='<' || c=='>' => format!("bitshift ('{c}{c}=')"),
Token::AssignmentOperation(c) => format!("operation ('{c}=')"),
Token::IncrementOperation => "increment operation".to_string(),
Token::DecrementOperation => "decrement operation".to_string(),
Token::Arrow => "->".to_string(),
Token::Unknown(c) => format!("unknown ('{c}')"),
Token::Trivia => "trivia".to_string(),
Token::End => "end".to_string(),
ExpectedToken::Token(token) => match token {
Token::Separator(c) => format!("'{c}'"),
Token::Paren(c) => format!("'{c}'"),
Token::Attribute => "@".to_string(),
Token::Number(_) => "number".to_string(),
Token::Word(s) => s.to_string(),
Token::Operation(c) => format!("operation ('{c}')"),
Token::LogicalOperation(c) => format!("logical operation ('{c}')"),
Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"),
Token::AssignmentOperation(c) if c == '<' || c == '>' => {
format!("bitshift ('{c}{c}=')")
}
}
Token::AssignmentOperation(c) => format!("operation ('{c}=')"),
Token::IncrementOperation => "increment operation".to_string(),
Token::DecrementOperation => "decrement operation".to_string(),
Token::Arrow => "->".to_string(),
Token::Unknown(c) => format!("unknown ('{c}')"),
Token::Trivia => "trivia".to_string(),
Token::End => "end".to_string(),
},
ExpectedToken::Identifier => "identifier".to_string(),
ExpectedToken::PrimaryExpression => "expression".to_string(),
ExpectedToken::Assignment => "assignment or increment/decrement".to_string(),
ExpectedToken::SwitchItem => "switch item ('case' or 'default') or a closing curly bracket to signify the end of the switch statement ('}')".to_string(),
ExpectedToken::WorkgroupSizeSeparator => "workgroup size separator (',') or a closing parenthesis".to_string(),
ExpectedToken::GlobalItem => "global item ('struct', 'const', 'var', 'alias', ';', 'fn') or the end of the file".to_string(),
ExpectedToken::SwitchItem => concat!(
"switch item ('case' or 'default') or a closing curly bracket ",
"to signify the end of the switch statement ('}')"
)
.to_string(),
ExpectedToken::WorkgroupSizeSeparator => {
"workgroup size separator (',') or a closing parenthesis".to_string()
}
ExpectedToken::GlobalItem => concat!(
"global item ('struct', 'const', 'var', 'alias', ';', 'fn') ",
"or the end of the file"
)
.to_string(),
ExpectedToken::Type => "type".to_string(),
ExpectedToken::Variable => "variable access".to_string(),
ExpectedToken::Function => "function name".to_string(),
@ -384,9 +394,11 @@ impl<'a> Error<'a> {
notes: vec![],
},
Error::BadIncrDecrReferenceType(span) => ParseError {
message:
"increment/decrement operation requires reference type to be one of i32 or u32"
.to_string(),
message: concat!(
"increment/decrement operation requires ",
"reference type to be one of i32 or u32"
)
.to_string(),
labels: vec![(span, "must be a reference type of i32 or u32".into())],
notes: vec![],
},
@ -527,25 +539,24 @@ impl<'a> Error<'a> {
labels: vec![(span, "type can't be inferred".into())],
notes: vec![],
},
Error::InitializationTypeMismatch { name, ref expected, ref got } => {
ParseError {
message: format!(
"the type of `{}` is expected to be `{}`, but got `{}`",
&source[name], expected, got,
),
labels: vec![(
name,
format!("definition of `{}`", &source[name]).into(),
)],
notes: vec![],
}
}
Error::InitializationTypeMismatch {
name,
ref expected,
ref got,
} => ParseError {
message: format!(
"the type of `{}` is expected to be `{}`, but got `{}`",
&source[name], expected, got,
),
labels: vec![(name, format!("definition of `{}`", &source[name]).into())],
notes: vec![],
},
Error::DeclMissingTypeAndInit(name_span) => ParseError {
message: format!("declaration of `{}` needs a type specifier or initializer", &source[name_span]),
labels: vec![(
name_span,
"needs a type specifier or initializer".into(),
)],
message: format!(
"declaration of `{}` needs a type specifier or initializer",
&source[name_span]
),
labels: vec![(name_span, "needs a type specifier or initializer".into())],
notes: vec![],
},
Error::MissingAttribute(name, name_span) => ParseError {
@ -725,7 +736,11 @@ impl<'a> Error<'a> {
notes: vec![message.into()],
},
Error::ExpectedConstExprConcreteIntegerScalar(span) => ParseError {
message: "must be a const-expression that resolves to a concrete integer scalar (u32 or i32)".to_string(),
message: concat!(
"must be a const-expression that ",
"resolves to a concrete integer scalar (u32 or i32)"
)
.to_string(),
labels: vec![(span, "must resolve to u32 or i32".into())],
notes: vec![],
},
@ -754,9 +769,17 @@ impl<'a> Error<'a> {
},
Error::AutoConversion(ref error) => {
// destructuring ensures all fields are handled
let AutoConversionError { dest_span, ref dest_type, source_span, ref source_type } = **error;
let AutoConversionError {
dest_span,
ref dest_type,
source_span,
ref source_type,
} = **error;
ParseError {
message: format!("automatic conversions cannot convert `{source_type}` to `{dest_type}`"),
message: format!(
"automatic conversions cannot convert `{}` to `{}`",
source_type, dest_type
),
labels: vec![
(
dest_span,
@ -765,72 +788,77 @@ impl<'a> Error<'a> {
(
source_span,
format!("this expression has type {source_type}").into(),
)
),
],
notes: vec![],
}
},
}
Error::AutoConversionLeafScalar(ref error) => {
let AutoConversionLeafScalarError { dest_span, ref dest_scalar, source_span, ref source_type } = **error;
let AutoConversionLeafScalarError {
dest_span,
ref dest_scalar,
source_span,
ref source_type,
} = **error;
ParseError {
message: format!("automatic conversions cannot convert elements of `{source_type}` to `{dest_scalar}`"),
message: format!(
"automatic conversions cannot convert elements of `{}` to `{}`",
source_type, dest_scalar
),
labels: vec![
(
dest_span,
format!("a value with elements of type {dest_scalar} is required here").into(),
format!(
"a value with elements of type {} is required here",
dest_scalar
)
.into(),
),
(
source_span,
format!("this expression has type {source_type}").into(),
)
),
],
notes: vec![],
}
},
}
Error::ConcretizationFailed(ref error) => {
let ConcretizationFailedError { expr_span, ref expr_type, ref scalar, ref inner } = **error;
let ConcretizationFailedError {
expr_span,
ref expr_type,
ref scalar,
ref inner,
} = **error;
ParseError {
message: format!("failed to convert expression to a concrete type: {}", inner),
labels: vec![
(
expr_span,
format!("this expression has type {}", expr_type).into(),
)
],
notes: vec![
format!("the expression should have been converted to have {} scalar type", scalar),
]
message: format!("failed to convert expression to a concrete type: {inner}"),
labels: vec![(
expr_span,
format!("this expression has type {expr_type}").into(),
)],
notes: vec![format!(
"the expression should have been converted to have {} scalar type",
scalar
)],
}
},
}
Error::ExceededLimitForNestedBraces { span, limit } => ParseError {
message: "brace nesting limit reached".into(),
labels: vec![(span, "limit reached at this brace".into())],
notes: vec![
format!("nesting limit is currently set to {limit}"),
],
notes: vec![format!("nesting limit is currently set to {limit}")],
},
Error::PipelineConstantIDValue(span) => ParseError {
message: "pipeline constant ID must be between 0 and 65535 inclusive".to_string(),
labels: vec![(
span,
"must be between 0 and 65535 inclusive".into(),
)],
labels: vec![(span, "must be between 0 and 65535 inclusive".into())],
notes: vec![],
},
Error::NotBool(span) => ParseError {
message: "must be a const-expression that resolves to a bool".to_string(),
labels: vec![(
span,
"must resolve to bool".into(),
)],
labels: vec![(span, "must resolve to bool".into())],
notes: vec![],
},
Error::ConstAssertFailed(span) => ParseError {
message: "const_assert failure".to_string(),
labels: vec![(
span,
"evaluates to false".into(),
)],
labels: vec![(span, "evaluates to false".into())],
notes: vec![],
},
}

6
naga/src/front/wgsl/lower/mod.rs
View File

@ -1778,12 +1778,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
return Ok(());
}
ast::StatementKind::Ignore(expr) => {
ast::StatementKind::Phony(expr) => {
let mut emitter = Emitter::default();
emitter.start(&ctx.function.expressions);
let _ = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
let value = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
block.extend(emitter.finish(&ctx.function.expressions));
ctx.named_expressions
.insert(value, ("phony".to_string(), stmt.span));
return Ok(());
}
};

2
naga/src/front/wgsl/parse/ast.rs
View File

@ -284,7 +284,7 @@ pub enum StatementKind<'a> {
},
Increment(Handle<Expression<'a>>),
Decrement(Handle<Expression<'a>>),
Ignore(Handle<Expression<'a>>),
Phony(Handle<Expression<'a>>),
ConstAssert(Handle<Expression<'a>>),
}

2
naga/src/front/wgsl/parse/mod.rs
View File

@ -1696,7 +1696,7 @@ impl Parser {
let expr = self.general_expression(lexer, ctx)?;
lexer.expect(Token::Separator(';'))?;
ast::StatementKind::Ignore(expr)
ast::StatementKind::Phony(expr)
}
"let" => {
let _ = lexer.next();

21
naga/src/lib.rs
View File

@ -1402,21 +1402,20 @@ pub enum Expression {
/// ## Dynamic indexing restrictions
///
/// To accommodate restrictions in some of the shader languages that Naga
/// targets, it is not permitted to subscript a matrix or array with a
/// dynamically computed index unless that matrix or array appears behind a
/// pointer. In other words, if the inner type of `base` is [`Array`] or
/// [`Matrix`], then `index` must be a constant. But if the type of `base`
/// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a
/// `size`, then the index may be any expression of integer type.
/// targets, it is not permitted to subscript a matrix with a dynamically
/// computed index unless that matrix appears behind a pointer. In other
/// words, if the inner type of `base` is [`Matrix`], then `index` must be a
/// constant. But if the type of `base` is a [`Pointer`] to an matrix, then
/// the index may be any expression of integer type.
///
/// You can use the [`Expression::is_dynamic_index`] method to determine
/// whether a given index expression requires matrix or array base operands
/// to be behind a pointer.
/// whether a given index expression requires matrix base operands to be
/// behind a pointer.
///
/// (It would be simpler to always require the use of `AccessIndex` when
/// subscripting arrays and matrices that are not behind pointers, but to
/// accommodate existing front ends, Naga also permits `Access`, with a
/// restricted `index`.)
/// subscripting matrices that are not behind pointers, but to accommodate
/// existing front ends, Naga also permits `Access`, with a restricted
/// `index`.)
///
/// [`Vector`]: TypeInner::Vector
/// [`Matrix`]: TypeInner::Matrix

8
naga/src/proc/mod.rs
View File

@ -521,12 +521,12 @@ impl crate::Expression {
}
}
/// Return true if this expression is a dynamic array index, for [`Access`].
/// Return true if this expression is a dynamic array/vector/matrix index,
/// for [`Access`].
///
/// This method returns true if this expression is a dynamically computed
/// index, and as such can only be used to index matrices and arrays when
/// they appear behind a pointer. See the documentation for [`Access`] for
/// details.
/// index, and as such can only be used to index matrices when they appear
/// behind a pointer. See the documentation for [`Access`] for details.
///
/// Note, this does not check the _type_ of the given expression. It's up to
/// the caller to establish that the `Access` expression is well-typed

290
naga/src/proc/typifier.rs
View File

@ -92,6 +92,13 @@ pub enum TypeResolution {
/// available in the associated arena. However, the `TypeInner` itself may
/// contain `Handle<Type>` values referring to types from the arena.
///
/// The inner type must only be one of the following variants:
/// - TypeInner::Pointer
/// - TypeInner::ValuePointer
/// - TypeInner::Matrix (generated by matrix multiplication)
/// - TypeInner::Vector
/// - TypeInner::Scalar
///
/// [`TypeInner`]: crate::TypeInner
Value(crate::TypeInner),
}
@ -631,41 +638,37 @@ impl<'a> ResolveContext<'a> {
use crate::MathFunction as Mf;
let res_arg = past(arg)?;
match fun {
// comparison
Mf::Abs |
Mf::Min |
Mf::Max |
Mf::Clamp |
Mf::Saturate |
// trigonometry
Mf::Cos |
Mf::Cosh |
Mf::Sin |
Mf::Sinh |
Mf::Tan |
Mf::Tanh |
Mf::Acos |
Mf::Asin |
Mf::Atan |
Mf::Atan2 |
Mf::Asinh |
Mf::Acosh |
Mf::Atanh |
Mf::Radians |
Mf::Degrees |
// decomposition
Mf::Ceil |
Mf::Floor |
Mf::Round |
Mf::Fract |
Mf::Trunc |
Mf::Ldexp |
// exponent
Mf::Exp |
Mf::Exp2 |
Mf::Log |
Mf::Log2 |
Mf::Pow => res_arg.clone(),
Mf::Abs
| Mf::Min
| Mf::Max
| Mf::Clamp
| Mf::Saturate
| Mf::Cos
| Mf::Cosh
| Mf::Sin
| Mf::Sinh
| Mf::Tan
| Mf::Tanh
| Mf::Acos
| Mf::Asin
| Mf::Atan
| Mf::Atan2
| Mf::Asinh
| Mf::Acosh
| Mf::Atanh
| Mf::Radians
| Mf::Degrees
| Mf::Ceil
| Mf::Floor
| Mf::Round
| Mf::Fract
| Mf::Trunc
| Mf::Ldexp
| Mf::Exp
| Mf::Exp2
| Mf::Log
| Mf::Log2
| Mf::Pow => res_arg.clone(),
Mf::Modf | Mf::Frexp => {
let (size, width) = match res_arg.inner_with(types) {
&Ti::Scalar(crate::Scalar {
@ -673,77 +676,81 @@ impl<'a> ResolveContext<'a> {
width,
}) => (None, width),
&Ti::Vector {
scalar: crate::Scalar {
kind: crate::ScalarKind::Float,
width,
},
scalar:
crate::Scalar {
kind: crate::ScalarKind::Float,
width,
},
size,
} => (Some(size), width),
ref other =>
return Err(ResolveError::IncompatibleOperands(format!("{fun:?}({other:?}, _)")))
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?}, _)"
)))
}
};
let result = self
.special_types
.predeclared_types
.get(&if fun == Mf::Modf {
crate::PredeclaredType::ModfResult { size, width }
} else {
crate::PredeclaredType::FrexpResult { size, width }
})
.ok_or(ResolveError::MissingSpecialType)?;
.special_types
.predeclared_types
.get(&if fun == Mf::Modf {
crate::PredeclaredType::ModfResult { size, width }
} else {
crate::PredeclaredType::FrexpResult { size, width }
})
.ok_or(ResolveError::MissingSpecialType)?;
TypeResolution::Handle(*result)
},
// geometry
}
Mf::Dot => match *res_arg.inner_with(types) {
Ti::Vector {
size: _,
scalar,
} => TypeResolution::Value(Ti::Scalar(scalar)),
ref other =>
return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?}, _)")
)),
Ti::Vector { size: _, scalar } => TypeResolution::Value(Ti::Scalar(scalar)),
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?}, _)"
)))
}
},
Mf::Outer => {
let arg1 = arg1.ok_or_else(|| ResolveError::IncompatibleOperands(
format!("{fun:?}(_, None)")
))?;
let arg1 = arg1.ok_or_else(|| {
ResolveError::IncompatibleOperands(format!("{fun:?}(_, None)"))
})?;
match (res_arg.inner_with(types), past(arg1)?.inner_with(types)) {
(
&Ti::Vector { size: columns, scalar },
&Ti::Vector{ size: rows, .. }
&Ti::Vector {
size: columns,
scalar,
},
&Ti::Vector { size: rows, .. },
) => TypeResolution::Value(Ti::Matrix {
columns,
rows,
scalar,
}),
(left, right) =>
return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({left:?}, {right:?})")
)),
(left, right) => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({left:?}, {right:?})"
)))
}
}
}
Mf::Cross => res_arg.clone(),
Mf::Distance | Mf::Length => match *res_arg.inner_with(types) {
Ti::Scalar(scalar) | Ti::Vector { scalar, size: _ } => {
TypeResolution::Value(Ti::Scalar(scalar))
}
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?})"
)))
}
},
Mf::Cross => res_arg.clone(),
Mf::Distance |
Mf::Length => match *res_arg.inner_with(types) {
Ti::Scalar(scalar) |
Ti::Vector {scalar,size:_} => TypeResolution::Value(Ti::Scalar(scalar)),
ref other => return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?})")
)),
},
Mf::Normalize |
Mf::FaceForward |
Mf::Reflect |
Mf::Refract => res_arg.clone(),
Mf::Normalize | Mf::FaceForward | Mf::Reflect | Mf::Refract => res_arg.clone(),
// computational
Mf::Sign |
Mf::Fma |
Mf::Mix |
Mf::Step |
Mf::SmoothStep |
Mf::Sqrt |
Mf::InverseSqrt => res_arg.clone(),
Mf::Sign
| Mf::Fma
| Mf::Mix
| Mf::Step
| Mf::SmoothStep
| Mf::Sqrt
| Mf::InverseSqrt => res_arg.clone(),
Mf::Transpose => match *res_arg.inner_with(types) {
Ti::Matrix {
columns,
@ -754,9 +761,11 @@ impl<'a> ResolveContext<'a> {
rows: columns,
scalar,
}),
ref other => return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?})")
)),
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?})"
)))
}
},
Mf::Inverse => match *res_arg.inner_with(types) {
Ti::Matrix {
@ -768,70 +777,75 @@ impl<'a> ResolveContext<'a> {
rows,
scalar,
}),
ref other => return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?})")
)),
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?})"
)))
}
},
Mf::Determinant => match *res_arg.inner_with(types) {
Ti::Matrix {
scalar,
..
} => TypeResolution::Value(Ti::Scalar(scalar)),
ref other => return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?})")
)),
Ti::Matrix { scalar, .. } => TypeResolution::Value(Ti::Scalar(scalar)),
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?})"
)))
}
},
// bits
Mf::CountTrailingZeros |
Mf::CountLeadingZeros |
Mf::CountOneBits |
Mf::ReverseBits |
Mf::ExtractBits |
Mf::InsertBits |
Mf::FirstTrailingBit |
Mf::FirstLeadingBit => match *res_arg.inner_with(types) {
Ti::Scalar(scalar @ crate::Scalar {
kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
..
}) => TypeResolution::Value(Ti::Scalar(scalar)),
Ti::Vector {
size,
scalar: scalar @ crate::Scalar {
Mf::CountTrailingZeros
| Mf::CountLeadingZeros
| Mf::CountOneBits
| Mf::ReverseBits
| Mf::ExtractBits
| Mf::InsertBits
| Mf::FirstTrailingBit
| Mf::FirstLeadingBit => match *res_arg.inner_with(types) {
Ti::Scalar(
scalar @ crate::Scalar {
kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
..
}
},
) => TypeResolution::Value(Ti::Scalar(scalar)),
Ti::Vector {
size,
scalar:
scalar @ crate::Scalar {
kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
..
},
} => TypeResolution::Value(Ti::Vector { size, scalar }),
ref other => return Err(ResolveError::IncompatibleOperands(
format!("{fun:?}({other:?})")
)),
ref other => {
return Err(ResolveError::IncompatibleOperands(format!(
"{fun:?}({other:?})"
)))
}
},
// data packing
Mf::Pack4x8snorm |
Mf::Pack4x8unorm |
Mf::Pack2x16snorm |
Mf::Pack2x16unorm |
Mf::Pack2x16float |
Mf::Pack4xI8 |
Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)),
Mf::Pack4x8snorm
| Mf::Pack4x8unorm
| Mf::Pack2x16snorm
| Mf::Pack2x16unorm
| Mf::Pack2x16float
| Mf::Pack4xI8
| Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)),
// data unpacking
Mf::Unpack4x8snorm |
Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector {
Mf::Unpack4x8snorm | Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector {
size: crate::VectorSize::Quad,
scalar: crate::Scalar::F32
}),
Mf::Unpack2x16snorm |
Mf::Unpack2x16unorm |
Mf::Unpack2x16float => TypeResolution::Value(Ti::Vector {
size: crate::VectorSize::Bi,
scalar: crate::Scalar::F32
scalar: crate::Scalar::F32,
}),
Mf::Unpack2x16snorm | Mf::Unpack2x16unorm | Mf::Unpack2x16float => {
TypeResolution::Value(Ti::Vector {
size: crate::VectorSize::Bi,
scalar: crate::Scalar::F32,
})
}
Mf::Unpack4xI8 => TypeResolution::Value(Ti::Vector {
size: crate::VectorSize::Quad,
scalar: crate::Scalar::I32
scalar: crate::Scalar::I32,
}),
Mf::Unpack4xU8 => TypeResolution::Value(Ti::Vector {
size: crate::VectorSize::Quad,
scalar: crate::Scalar::U32
scalar: crate::Scalar::U32,
}),
}
}

1
naga/src/span.rs
View File

@ -11,6 +11,7 @@ pub struct Span {
impl Span {
pub const UNDEFINED: Self = Self { start: 0, end: 0 };
/// Creates a new `Span` from a range of byte indices
///
/// Note: end is exclusive, it doesn't belong to the `Span`

15
naga/src/valid/analyzer.rs
View File

@ -589,23 +589,16 @@ impl FunctionInfo {
requirements: UniformityRequirements::empty(),
}
}
// depends on the builtin or interpolation
// depends on the builtin
E::FunctionArgument(index) => {
let arg = &resolve_context.arguments[index as usize];
let uniform = match arg.binding {
Some(crate::Binding::BuiltIn(
// per-polygon built-ins are uniform
crate::BuiltIn::FrontFacing
// per-work-group built-ins are uniform
| crate::BuiltIn::WorkGroupId
crate::BuiltIn::WorkGroupId
| crate::BuiltIn::WorkGroupSize
| crate::BuiltIn::NumWorkGroups)
) => true,
// only flat inputs are uniform
Some(crate::Binding::Location {
interpolation: Some(crate::Interpolation::Flat),
..
}) => true,
| crate::BuiltIn::NumWorkGroups,
)) => true,
_ => false,
};
Uniformity {

7
naga/src/valid/expression.rs
View File

@ -240,9 +240,10 @@ impl super::Validator {
let base_type = &resolver[base];
// See the documentation for `Expression::Access`.
let dynamic_indexing_restricted = match *base_type {
Ti::Vector { .. } => false,
Ti::Matrix { .. } | Ti::Array { .. } => true,
Ti::Pointer { .. }
Ti::Matrix { .. } => true,
Ti::Vector { .. }
| Ti::Array { .. }
| Ti::Pointer { .. }
| Ti::ValuePointer { size: Some(_), .. }
| Ti::BindingArray { .. } => false,
ref other => {

4
naga/src/valid/type.rs
View File

@ -664,9 +664,6 @@ impl super::Validator {
)
}
Ti::BindingArray { base, size } => {
if base >= handle {
return Err(TypeError::InvalidArrayBaseType(base));
}
let type_info_mask = match size {
crate::ArraySize::Constant(_) => TypeFlags::SIZED | TypeFlags::HOST_SHAREABLE,
crate::ArraySize::Dynamic => {
@ -680,7 +677,6 @@ impl super::Validator {
// Currently Naga only supports binding arrays of structs for non-handle types.
match gctx.types[base].inner {
crate::TypeInner::Struct { .. } => {}
crate::TypeInner::Array { .. } => {}
_ => return Err(TypeError::BindingArrayBaseTypeNotStruct(base)),
};
}

2
naga/tests/in/6220-break-from-loop.param.ron Normal file
View File

@ -0,0 +1,2 @@
(
)

43
naga/tests/in/6220-break-from-loop.wgsl Normal file
View File

@ -0,0 +1,43 @@
// #6220: Don't generate unreachable SPIR-V blocks that branch into
// structured control flow constructs.
//
// Suppose we have Naga code like this:
//
// Block {
// ... prelude
// Block { ... nested }
// ... postlude
// }
//
// The SPIR-V back end used to always generate three separate SPIR-V
// blocks for the sections labeled "prelude", "nested", and
// "postlude", each block ending with a branch to the next, even if
// they were empty.
//
// However, the function below generates code that includes the
// following structure:
//
// Loop {
// body: Block {
// ... prelude
// Block { Break }
// ... postlude
// }
// continuing: ...
// }
//
// In this case, even though the `Break` renders the "postlude"
// unreachable, we used to generate a SPIR-V block for it anyway,
// ending with a branch to the `Loop`'s "continuing" block. However,
// SPIR-V's structured control flow rules forbid branches to a loop
// construct's continue target from outside the loop, so the SPIR-V
// module containing the unreachable block didn't pass validation.
//
// One might assume that unreachable blocks shouldn't affect
// validation, but the spec doesn't clearly agree, and this doesn't
// seem to be the way validation has been implemented.
fn break_from_loop() {
for (var i = 0; i < 4; i += 1) {
break;
}
}

11
naga/tests/in/access.wgsl
View File

@ -167,3 +167,14 @@ fn assign_through_ptr() {
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
assign_array_through_ptr_fn(&arr);
}
@vertex
fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let arr = array<i32, 5>(1, 2, 3, 4, 5);
let value = arr[vi];
return vec4<f32>(vec4<i32>(value));
}
fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
return a[i];
}

2
naga/tests/in/phony_assignment.ron Normal file
View File

@ -0,0 +1,2 @@
(
)

18
naga/tests/in/phony_assignment.wgsl Normal file
View File

@ -0,0 +1,18 @@
@group(0) @binding(0) var<uniform> binding: f32;
fn five() -> i32 {
return 5;
}
@compute @workgroup_size(1) fn main(
@builtin(global_invocation_id) id: vec3<u32>
) {
_ = binding;
_ = binding;
let a = 5;
_ = a;
_ = five();
let b = five();
// check for name collision
let phony = binding;
}

BIN
naga/tests/in/spv/atomic_exchange.spv Normal file
View File

Binary file not shown.

88
naga/tests/in/spv/atomic_exchange.spvasm Normal file
View File

@ -0,0 +1,88 @@
; SPIR-V
; Version: 1.5
; Generator: Google rspirv; 0
; Bound: 63
; Schema: 0
OpCapability Shader
OpCapability VulkanMemoryModel
OpMemoryModel Logical Vulkan
OpEntryPoint GLCompute %1 "stage::test_atomic_exchange" %2 %3
OpExecutionMode %1 LocalSize 32 1 1
OpMemberDecorate %_struct_11 0 Offset 0
OpMemberDecorate %_struct_11 1 Offset 4
OpDecorate %_struct_12 Block
OpMemberDecorate %_struct_12 0 Offset 0
OpDecorate %2 Binding 0
OpDecorate %2 DescriptorSet 0
OpDecorate %3 NonWritable
OpDecorate %3 Binding 1
OpDecorate %3 DescriptorSet 0
%uint = OpTypeInt 32 0
%void = OpTypeVoid
%15 = OpTypeFunction %void
%bool = OpTypeBool
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%false = OpConstantFalse %bool
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%uint_1 = OpConstant %uint 1
%_struct_11 = OpTypeStruct %uint %uint
%22 = OpUndef %_struct_11
%int = OpTypeInt 32 1
%true = OpConstantTrue %bool
%_struct_12 = OpTypeStruct %uint
%_ptr_StorageBuffer__struct_12 = OpTypePointer StorageBuffer %_struct_12
%2 = OpVariable %_ptr_StorageBuffer__struct_12 StorageBuffer
%3 = OpVariable %_ptr_StorageBuffer__struct_12 StorageBuffer
%26 = OpUndef %uint
%1 = OpFunction %void None %15
%27 = OpLabel
%28 = OpAccessChain %_ptr_StorageBuffer_uint %2 %uint_0
%29 = OpAccessChain %_ptr_StorageBuffer_uint %3 %uint_0
%30 = OpLoad %uint %29
%31 = OpCompositeConstruct %_struct_11 %uint_0 %30
OpBranch %32
%32 = OpLabel
%33 = OpPhi %_struct_11 %31 %27 %34 %35
%36 = OpPhi %uint %uint_0 %27 %37 %35
OpLoopMerge %38 %35 None
OpBranch %39
%39 = OpLabel
%40 = OpCompositeExtract %uint %33 0
%41 = OpCompositeExtract %uint %33 1
%42 = OpULessThan %bool %40 %41
OpSelectionMerge %43 None
OpBranchConditional %42 %44 %45
%44 = OpLabel
%47 = OpIAdd %uint %40 %uint_1
%49 = OpCompositeInsert %_struct_11 %47 %33 0
%50 = OpCompositeConstruct %_struct_11 %uint_1 %40
OpBranch %43
%45 = OpLabel
%51 = OpCompositeInsert %_struct_11 %uint_0 %22 0
OpBranch %43
%43 = OpLabel
%52 = OpPhi %_struct_11 %49 %44 %33 %45
%53 = OpPhi %_struct_11 %50 %44 %51 %45
%54 = OpCompositeExtract %uint %53 0
%55 = OpBitcast %int %54
OpSelectionMerge %56 None
OpSwitch %55 %57 0 %58 1 %59
%57 = OpLabel
OpBranch %56
%58 = OpLabel
OpBranch %56
%59 = OpLabel
%60 = OpAtomicExchange %uint %28 %uint_2 %uint_0 %36
%61 = OpIAdd %uint %36 %60
OpBranch %56
%56 = OpLabel
%62 = OpPhi %bool %false %57 %false %58 %true %59
%34 = OpPhi %_struct_11 %22 %57 %22 %58 %52 %59
%37 = OpPhi %uint %26 %57 %26 %58 %61 %59
OpBranch %35
%35 = OpLabel
OpBranchConditional %62 %32 %38
%38 = OpLabel
OpReturn
OpFunctionEnd

BIN
naga/tests/in/spv/atomic_i_add_sub.spv Normal file
View File

Binary file not shown.

51
naga/tests/in/spv/atomic_i_add_sub.spvasm Normal file
View File

@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.5
; Generator: Google rspirv; 0
; Bound: 30
; Schema: 0
OpCapability Shader
OpCapability VulkanMemoryModel
OpMemoryModel Logical Vulkan
OpEntryPoint GLCompute %1 "stage::test_atomic_i_add_sub" %2 %3
OpExecutionMode %1 LocalSize 32 1 1
OpDecorate %_runtimearr_uint ArrayStride 4
OpDecorate %_struct_7 Block
OpMemberDecorate %_struct_7 0 Offset 0
OpDecorate %_struct_8 Block
OpMemberDecorate %_struct_8 0 Offset 0
OpDecorate %2 Binding 0
OpDecorate %2 DescriptorSet 0
OpDecorate %3 Binding 1
OpDecorate %3 DescriptorSet 0
%uint = OpTypeInt 32 0
%void = OpTypeVoid
%11 = OpTypeFunction %void
%bool = OpTypeBool
%_runtimearr_uint = OpTypeRuntimeArray %uint
%_struct_7 = OpTypeStruct %_runtimearr_uint
%_ptr_StorageBuffer__struct_7 = OpTypePointer StorageBuffer %_struct_7
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%_struct_8 = OpTypeStruct %uint
%_ptr_StorageBuffer__struct_8 = OpTypePointer StorageBuffer %_struct_8
%2 = OpVariable %_ptr_StorageBuffer__struct_8 StorageBuffer
%3 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
%1 = OpFunction %void None %11
%19 = OpLabel
%20 = OpAccessChain %_ptr_StorageBuffer_uint %2 %uint_0
%22 = OpArrayLength %uint %3 0
%23 = OpAtomicIAdd %uint %20 %uint_2 %uint_0 %uint_2
%24 = OpAtomicISub %uint %20 %uint_2 %uint_0 %23
%25 = OpULessThan %bool %23 %22
OpSelectionMerge %26 None
OpBranchConditional %25 %27 %28
%27 = OpLabel
%29 = OpAccessChain %_ptr_StorageBuffer_uint %3 %uint_0 %23
OpStore %29 %24
OpBranch %26
%28 = OpLabel
OpBranch %26
%26 = OpLabel
OpReturn
OpFunctionEnd

BIN
naga/tests/in/spv/atomic_i_decrement.spv Normal file
View File

Binary file not shown.

64
naga/tests/in/spv/atomic_i_decrement.spvasm Normal file
View File

@ -0,0 +1,64 @@
; SPIR-V
; Version: 1.5
; Generator: Google rspirv; 0
; Bound: 42
; Schema: 0
OpCapability Shader
OpCapability VulkanMemoryModel
OpMemoryModel Logical Vulkan
OpEntryPoint GLCompute %1 "stage::test_atomic_i_decrement" %2 %3
OpExecutionMode %1 LocalSize 32 1 1
OpDecorate %_runtimearr_uint ArrayStride 4
OpDecorate %_struct_7 Block
OpMemberDecorate %_struct_7 0 Offset 0
OpDecorate %_struct_8 Block
OpMemberDecorate %_struct_8 0 Offset 0
OpDecorate %2 Binding 0
OpDecorate %2 DescriptorSet 0
OpDecorate %3 Binding 1
OpDecorate %3 DescriptorSet 0
%uint = OpTypeInt 32 0
%void = OpTypeVoid
%11 = OpTypeFunction %void
%bool = OpTypeBool
%_runtimearr_uint = OpTypeRuntimeArray %uint
%_struct_7 = OpTypeStruct %_runtimearr_uint
%_ptr_StorageBuffer__struct_7 = OpTypePointer StorageBuffer %_struct_7
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%false = OpConstantFalse %bool
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%true = OpConstantTrue %bool
%_struct_8 = OpTypeStruct %uint
%_ptr_StorageBuffer__struct_8 = OpTypePointer StorageBuffer %_struct_8
%2 = OpVariable %_ptr_StorageBuffer__struct_8 StorageBuffer
%3 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
%1 = OpFunction %void None %11
%21 = OpLabel
%22 = OpAccessChain %_ptr_StorageBuffer_uint %2 %uint_0
%24 = OpArrayLength %uint %3 0
OpBranch %25
%25 = OpLabel
OpLoopMerge %26 %27 None
OpBranch %28
%28 = OpLabel
%29 = OpAtomicIDecrement %uint %22 %uint_2 %uint_0
%30 = OpULessThan %bool %29 %24
OpSelectionMerge %31 None
OpBranchConditional %30 %32 %33
%32 = OpLabel
%34 = OpAccessChain %_ptr_StorageBuffer_uint %3 %uint_0 %29
OpStore %34 %29
%35 = OpIEqual %bool %29 %uint_0
%41 = OpSelect %bool %35 %false %true
OpBranch %31
%33 = OpLabel
OpBranch %31
%31 = OpLabel
%40 = OpPhi %bool %41 %32 %false %33
OpBranch %27
%27 = OpLabel
OpBranchConditional %40 %25 %26
%26 = OpLabel
OpReturn
OpFunctionEnd

1
naga/tests/in/spv/atomic_i_increment.spvasm
View File

@ -59,4 +59,3 @@
%26 = OpLabel
OpReturn
OpFunctionEnd

BIN
naga/tests/in/spv/atomic_load_and_store.spv Normal file
View File

Binary file not shown.

86
naga/tests/in/spv/atomic_load_and_store.spvasm Normal file
View File

@ -0,0 +1,86 @@
; SPIR-V
; Version: 1.5
; Generator: Google rspirv; 0
; Bound: 60
; Schema: 0
OpCapability Shader
OpCapability VulkanMemoryModel
OpMemoryModel Logical Vulkan
OpEntryPoint GLCompute %1 "stage::test_atomic_load_and_store" %2 %3
OpExecutionMode %1 LocalSize 32 1 1
OpMemberDecorate %_struct_11 0 Offset 0
OpMemberDecorate %_struct_11 1 Offset 4
OpDecorate %_struct_12 Block
OpMemberDecorate %_struct_12 0 Offset 0
OpDecorate %2 Binding 0
OpDecorate %2 DescriptorSet 0
OpDecorate %3 NonWritable
OpDecorate %3 Binding 1
OpDecorate %3 DescriptorSet 0
%uint = OpTypeInt 32 0
%void = OpTypeVoid
%15 = OpTypeFunction %void
%bool = OpTypeBool
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%false = OpConstantFalse %bool
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%uint_1 = OpConstant %uint 1
%_struct_11 = OpTypeStruct %uint %uint
%22 = OpUndef %_struct_11
%int = OpTypeInt 32 1
%true = OpConstantTrue %bool
%_struct_12 = OpTypeStruct %uint
%_ptr_StorageBuffer__struct_12 = OpTypePointer StorageBuffer %_struct_12
%2 = OpVariable %_ptr_StorageBuffer__struct_12 StorageBuffer
%3 = OpVariable %_ptr_StorageBuffer__struct_12 StorageBuffer
%1 = OpFunction %void None %15
%26 = OpLabel
%27 = OpAccessChain %_ptr_StorageBuffer_uint %2 %uint_0
%28 = OpAccessChain %_ptr_StorageBuffer_uint %3 %uint_0
%29 = OpLoad %uint %28
%30 = OpCompositeConstruct %_struct_11 %uint_0 %29
OpBranch %31
%31 = OpLabel
%32 = OpPhi %_struct_11 %30 %26 %33 %34
OpLoopMerge %35 %34 None
OpBranch %36
%36 = OpLabel
%37 = OpCompositeExtract %uint %32 0
%38 = OpCompositeExtract %uint %32 1
%39 = OpULessThan %bool %37 %38
OpSelectionMerge %40 None
OpBranchConditional %39 %41 %42
%41 = OpLabel
%44 = OpIAdd %uint %37 %uint_1
%46 = OpCompositeInsert %_struct_11 %44 %32 0
%47 = OpCompositeConstruct %_struct_11 %uint_1 %37
OpBranch %40
%42 = OpLabel
%48 = OpCompositeInsert %_struct_11 %uint_0 %22 0
OpBranch %40
%40 = OpLabel
%49 = OpPhi %_struct_11 %46 %41 %32 %42
%50 = OpPhi %_struct_11 %47 %41 %48 %42
%51 = OpCompositeExtract %uint %50 0
%52 = OpBitcast %int %51
OpSelectionMerge %53 None
OpSwitch %52 %54 0 %55 1 %56
%54 = OpLabel
OpBranch %53
%55 = OpLabel
OpBranch %53
%56 = OpLabel
%57 = OpAtomicLoad %uint %27 %uint_2 %uint_0
%58 = OpIAdd %uint %57 %uint_2
OpAtomicStore %27 %uint_2 %uint_0 %58
OpBranch %53
%53 = OpLabel
%59 = OpPhi %bool %false %54 %false %55 %true %56
%33 = OpPhi %_struct_11 %22 %54 %22 %55 %49 %56
OpBranch %34
%34 = OpLabel
OpBranchConditional %59 %31 %35
%35 = OpLabel
OpReturn
OpFunctionEnd

186
naga/tests/out/analysis/access.info.ron
View File

@ -2735,6 +2735,54 @@
sampling: [],
dual_source_blending: false,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
may_kill: false,
sampling_set: [],
global_uses: [
(""),
(""),
(""),
(""),
(""),
],
expressions: [
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(25),
),
(
uniformity: (
non_uniform_result: Some(1),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
],
sampling: [],
dual_source_blending: false,
),
],
entry_points: [
(
@ -3981,6 +4029,144 @@
sampling: [],
dual_source_blending: false,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
may_kill: false,
sampling_set: [],
global_uses: [
(""),
(""),
(""),
(""),
(""),
],
expressions: [
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(0),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Sint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Sint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Sint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Sint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Scalar((
kind: Sint,
width: 4,
))),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(25),
),
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Sint,
width: 4,
),
)),
),
(
uniformity: (
non_uniform_result: Some(0),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Value(Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
)),
),
],
sampling: [],
dual_source_blending: false,
),
],
const_expression_types: [
Value(Scalar((

12
naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
View File

@ -20,8 +20,8 @@ struct MatCx2InArray {
mat4x2 am[2];
};
float read_from_private(inout float foo_1) {
float _e1 = foo_1;
float read_from_private(inout float foo_2) {
float _e1 = foo_2;
return _e1;
}
@ -34,11 +34,15 @@ void assign_through_ptr_fn(inout uint p) {
return;
}
void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
foo_3 = vec4[2](vec4(1.0), vec4(2.0));
return;
}
int array_by_value(int a_1[5], int i) {
return a_1[i];
}
void main() {
uint val = 33u;
vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));

52
naga/tests/out/glsl/access.foo.Vertex.glsl Normal file
View File

@ -0,0 +1,52 @@
#version 310 es
precision highp float;
precision highp int;
struct GlobalConst {
uint a;
uvec3 b;
int c;
};
struct AlignedWrapper {
int value;
};
struct Baz {
mat3x2 m;
};
struct MatCx2InArray {
mat4x2 am[2];
};
float read_from_private(inout float foo_2) {
float _e1 = foo_2;
return _e1;
}
float test_arr_as_arg(float a[5][10]) {
return a[4][9];
}
void assign_through_ptr_fn(inout uint p) {
p = 42u;
return;
}
void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
foo_3 = vec4[2](vec4(1.0), vec4(2.0));
return;
}
int array_by_value(int a_1[5], int i) {
return a_1[i];
}
void main() {
uint vi_1 = uint(gl_VertexID);
int arr_1[5] = int[5](1, 2, 3, 4, 5);
int value = arr_1[vi_1];
gl_Position = vec4(ivec4(value));
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;
}

12
naga/tests/out/glsl/access.foo_frag.Fragment.glsl
View File

@ -30,8 +30,8 @@ layout(std430) buffer type_13_block_1Fragment { ivec2 _group_0_binding_2_fs; };
layout(location = 0) out vec4 _fs2p_location0;
float read_from_private(inout float foo_1) {
float _e1 = foo_1;
float read_from_private(inout float foo_2) {
float _e1 = foo_2;
return _e1;
}
@ -44,11 +44,15 @@ void assign_through_ptr_fn(inout uint p) {
return;
}
void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
foo_3 = vec4[2](vec4(1.0), vec4(2.0));
return;
}
int array_by_value(int a_1[5], int i) {
return a_1[i];
}
void main() {
_group_0_binding_0_fs._matrix[1][2] = 1.0;
_group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));

16
naga/tests/out/glsl/access.foo_vert.Vertex.glsl
View File

@ -103,8 +103,8 @@ void test_matrix_within_array_within_struct_accesses() {
return;
}
float read_from_private(inout float foo_1) {
float _e1 = foo_1;
float read_from_private(inout float foo_2) {
float _e1 = foo_2;
return _e1;
}
@ -117,11 +117,15 @@ void assign_through_ptr_fn(inout uint p) {
return;
}
void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
foo_2 = vec4[2](vec4(1.0), vec4(2.0));
void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
foo_3 = vec4[2](vec4(1.0), vec4(2.0));
return;
}
int array_by_value(int a_1[5], int i) {
return a_1[i];
}
void main() {
uint vi = uint(gl_VertexID);
float foo = 0.0;
@ -133,10 +137,10 @@ void main() {
mat4x3 _matrix = _group_0_binding_0_vs._matrix;
uvec2 arr_1[2] = _group_0_binding_0_vs.arr;
float b = _group_0_binding_0_vs._matrix[3u][0];
int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
ivec2 c = _group_0_binding_2_vs;
float _e33 = read_from_private(foo);
c2_ = int[5](a_1, int(b), 3, 4, 5);
c2_ = int[5](a_2, int(b), 3, 4, 5);
c2_[(vi + 1u)] = 42;
int value = c2_[vi];
float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));

23
naga/tests/out/glsl/phony_assignment.main.Compute.glsl Normal file
View File

@ -0,0 +1,23 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
uniform type_block_0Compute { float _group_0_binding_0_cs; };
int five() {
return 5;
}
void main() {
uvec3 id = gl_GlobalInvocationID;
float phony = _group_0_binding_0_cs;
float phony_1 = _group_0_binding_0_cs;
int _e6 = five();
int _e7 = five();
float phony_2 = _group_0_binding_0_cs;
}

24
naga/tests/out/hlsl/access.hlsl
View File

@ -201,9 +201,9 @@ void test_matrix_within_array_within_struct_accesses()
return;
}
float read_from_private(inout float foo_1)
float read_from_private(inout float foo_2)
{
float _e1 = foo_1;
float _e1 = foo_2;
return _e1;
}
@ -224,12 +224,17 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) {
return ret;
}
void assign_array_through_ptr_fn(inout float4 foo_2[2])
void assign_array_through_ptr_fn(inout float4 foo_3[2])
{
foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
foo_3 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
return;
}
int array_by_value(int a_1[5], int i)
{
return a_1[i];
}
typedef int ret_Constructarray5_int_[5];
ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) {
int ret[5] = { arg0, arg1, arg2, arg3, arg4 };
@ -266,10 +271,10 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8)));
float b = asfloat(bar.Load(0+3u*16+0));
int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
int2 c = asint(qux.Load2(0));
const float _e33 = read_from_private(foo);
c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5);
c2_ = Constructarray5_int_(a_2, int(b), 3, 4, 5);
c2_[(vi + 1u)] = 42;
int value = c2_[vi];
const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__());
@ -310,3 +315,10 @@ void assign_through_ptr()
assign_array_through_ptr_fn(arr);
return;
}
float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position
{
int arr_2[5] = Constructarray5_int_(1, 2, 3, 4, 5);
int value_1 = arr_2[vi_1];
return float4((value_1).xxxx);
}

4
naga/tests/out/hlsl/access.ron
View File

@ -4,6 +4,10 @@
entry_point:"foo_vert",
target_profile:"vs_5_1",
),
(
entry_point:"foo_1",
target_profile:"vs_5_1",
),
],
fragment:[
(

16
naga/tests/out/hlsl/phony_assignment.hlsl Normal file
View File

@ -0,0 +1,16 @@
cbuffer binding : register(b0) { float binding; }
int five()
{
return 5;
}
[numthreads(1, 1, 1)]
void main(uint3 id : SV_DispatchThreadID)
{
float phony = binding;
float phony_1 = binding;
const int _e6 = five();
const int _e7 = five();
float phony_2 = binding;
}

12
naga/tests/out/hlsl/phony_assignment.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"main",
target_profile:"cs_5_1",
),
],
)

117
naga/tests/out/ir/access.compact.ron
View File

@ -1655,6 +1655,47 @@
),
],
),
(
name: Some("array_by_value"),
arguments: [
(
name: Some("a"),
ty: 25,
binding: None,
),
(
name: Some("i"),
ty: 2,
binding: None,
),
],
result: Some((
ty: 2,
binding: None,
)),
local_variables: [],
expressions: [
FunctionArgument(0),
FunctionArgument(1),
Access(
base: 0,
index: 1,
),
],
named_expressions: {
0: "a",
1: "i",
},
body: [
Emit((
start: 2,
end: 3,
)),
Return(
value: Some(2),
),
],
),
],
entry_points: [
(
@ -2230,5 +2271,81 @@
],
),
),
(
name: "foo",
stage: Vertex,
early_depth_test: None,
workgroup_size: (0, 0, 0),
function: (
name: Some("foo"),
arguments: [
(
name: Some("vi"),
ty: 0,
binding: Some(BuiltIn(VertexIndex)),
),
],
result: Some((
ty: 24,
binding: Some(BuiltIn(Position(
invariant: false,
))),
)),
local_variables: [],
expressions: [
FunctionArgument(0),
Literal(I32(1)),
Literal(I32(2)),
Literal(I32(3)),
Literal(I32(4)),
Literal(I32(5)),
Compose(
ty: 25,
components: [
1,
2,
3,
4,
5,
],
),
Access(
base: 6,
index: 0,
),
Splat(
size: Quad,
value: 7,
),
As(
expr: 8,
kind: Float,
convert: Some(4),
),
],
named_expressions: {
0: "vi",
6: "arr",
7: "value",
},
body: [
Emit((
start: 6,
end: 7,
)),
Emit((
start: 7,
end: 8,
)),
Emit((
start: 8,
end: 10,
)),
Return(
value: Some(9),
),
],
),
),
],
)

117
naga/tests/out/ir/access.ron
View File

@ -1655,6 +1655,47 @@
),
],
),
(
name: Some("array_by_value"),
arguments: [
(
name: Some("a"),
ty: 25,
binding: None,
),
(
name: Some("i"),
ty: 2,
binding: None,
),
],
result: Some((
ty: 2,
binding: None,
)),
local_variables: [],
expressions: [
FunctionArgument(0),
FunctionArgument(1),
Access(
base: 0,
index: 1,
),
],
named_expressions: {
0: "a",
1: "i",
},
body: [
Emit((
start: 2,
end: 3,
)),
Return(
value: Some(2),
),
],
),
],
entry_points: [
(
@ -2230,5 +2271,81 @@
],
),
),
(
name: "foo",
stage: Vertex,
early_depth_test: None,
workgroup_size: (0, 0, 0),
function: (
name: Some("foo"),
arguments: [
(
name: Some("vi"),
ty: 0,
binding: Some(BuiltIn(VertexIndex)),
),
],
result: Some((
ty: 24,
binding: Some(BuiltIn(Position(
invariant: false,
))),
)),
local_variables: [],
expressions: [
FunctionArgument(0),
Literal(I32(1)),
Literal(I32(2)),
Literal(I32(3)),
Literal(I32(4)),
Literal(I32(5)),
Compose(
ty: 25,
components: [
1,
2,
3,
4,
5,
],
),
Access(
base: 6,
index: 0,
),
Splat(
size: Quad,
value: 7,
),
As(
expr: 8,
kind: Float,
convert: Some(4),
),
],
named_expressions: {
0: "vi",
6: "arr",
7: "value",
},
body: [
Emit((
start: 6,
end: 7,
)),
Emit((
start: 7,
end: 8,
)),
Emit((
start: 8,
end: 10,
)),
Return(
value: Some(9),
),
],
),
),
],
)

4
naga/tests/out/ir/atomic_i_increment.compact.ron
View File

@ -216,10 +216,6 @@
),
],
reject: [
Emit((
start: 13,
end: 14,
)),
Atomic(
pointer: 7,
fun: Add,

4
naga/tests/out/ir/atomic_i_increment.ron
View File

@ -241,10 +241,6 @@
),
],
reject: [
Emit((
start: 14,
end: 15,
)),
Atomic(
pointer: 8,
fun: Add,

33
naga/tests/out/msl/access.msl
View File

@ -133,9 +133,9 @@ void test_matrix_within_array_within_struct_accesses(
}
float read_from_private(
thread float& foo_1
thread float& foo_2
) {
float _e1 = foo_1;
float _e1 = foo_2;
return _e1;
}
@ -153,12 +153,19 @@ void assign_through_ptr_fn(
}
void assign_array_through_ptr_fn(
thread type_22& foo_2
thread type_22& foo_3
) {
foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)};
foo_3 = type_22 {metal::float4(1.0), metal::float4(2.0)};
return;
}
int array_by_value(
type_20 a_1,
int i
) {
return a_1.inner[i];
}
struct foo_vertInput {
};
struct foo_vertOutput {
@ -181,10 +188,10 @@ vertex foo_vertOutput foo_vert(
metal::float4x3 _matrix = bar._matrix;
type_10 arr_1 = bar.arr;
float b = bar._matrix[3u].x;
int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
metal::int2 c = qux;
float _e33 = read_from_private(foo);
c2_ = type_20 {a_1, static_cast<int>(b), 3, 4, 5};
c2_ = type_20 {a_2, static_cast<int>(b), 3, 4, 5};
c2_.inner[vi + 1u] = 42;
int value = c2_.inner[vi];
float _e47 = test_arr_as_arg(type_18 {});
@ -217,3 +224,17 @@ kernel void assign_through_ptr(
assign_array_through_ptr_fn(arr);
return;
}
struct foo_1Input {
};
struct foo_1Output {
metal::float4 member_3 [[position]];
};
vertex foo_1Output foo_1(
uint vi_1 [[vertex_id]]
) {
type_20 arr_2 = type_20 {1, 2, 3, 4, 5};
int value_1 = arr_2.inner[vi_1];
return foo_1Output { static_cast<metal::float4>(metal::int4(value_1)) };
}

3
naga/tests/out/msl/boids.msl
View File

@ -55,8 +55,9 @@ kernel void main_(
vPos = _e8;
metal::float2 _e14 = particlesSrc.particles[index].vel;
vVel = _e14;
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
bool loop_init = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init) {
uint _e91 = i;
i = _e91 + 1u;

9
naga/tests/out/msl/break-if.msl
View File

@ -7,8 +7,9 @@ using metal::uint;
void breakIfEmpty(
) {
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
bool loop_init = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init) {
if (true) {
break;
@ -25,7 +26,7 @@ void breakIfEmptyBody(
bool b = {};
bool c = {};
bool loop_init_1 = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init_1) {
b = a;
bool _e2 = b;
@ -46,7 +47,7 @@ void breakIf(
bool d = {};
bool e = {};
bool loop_init_2 = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init_2) {
bool _e5 = e;
if (a_1 == e) {
@ -65,7 +66,7 @@ void breakIfSeparateVariable(
) {
uint counter = 0u;
bool loop_init_3 = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init_3) {
uint _e5 = counter;
if (counter == 5u) {

3
naga/tests/out/msl/collatz.msl
View File

@ -19,7 +19,8 @@ uint collatz_iterations(
uint n = {};
uint i = 0u;
n = n_base;
while(true) {
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
LOOP_IS_REACHABLE while(true) {
uint _e4 = n;
if (_e4 > 1u) {
} else {

13
naga/tests/out/msl/control-flow.msl
View File

@ -31,7 +31,8 @@ void switch_case_break(
void loop_switch_continue(
int x
) {
while(true) {
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
LOOP_IS_REACHABLE while(true) {
switch(x) {
case 1: {
continue;
@ -49,7 +50,7 @@ void loop_switch_continue_nesting(
int y,
int z
) {
while(true) {
LOOP_IS_REACHABLE while(true) {
switch(x_1) {
case 1: {
continue;
@ -60,7 +61,7 @@ void loop_switch_continue_nesting(
continue;
}
default: {
while(true) {
LOOP_IS_REACHABLE while(true) {
switch(z) {
case 1: {
continue;
@ -85,7 +86,7 @@ void loop_switch_continue_nesting(
}
}
}
while(true) {
LOOP_IS_REACHABLE while(true) {
switch(y) {
case 1:
default: {
@ -108,7 +109,7 @@ void loop_switch_omit_continue_variable_checks(
int w
) {
int pos_1 = 0;
while(true) {
LOOP_IS_REACHABLE while(true) {
switch(x_2) {
case 1: {
pos_1 = 1;
@ -119,7 +120,7 @@ void loop_switch_omit_continue_variable_checks(
}
}
}
while(true) {
LOOP_IS_REACHABLE while(true) {
switch(x_2) {
case 1: {
break;

3
naga/tests/out/msl/do-while.msl
View File

@ -8,8 +8,9 @@ using metal::uint;
void fb1_(
thread bool& cond
) {
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
bool loop_init = true;
while(true) {
LOOP_IS_REACHABLE while(true) {
if (!loop_init) {
bool _e1 = cond;
if (!(cond)) {

Some files were not shown because too many files have changed in this diff Show More