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 remote-tracking branch 'origin/trunk' into pr/shader-flt32-atomic

Browse Source
This commit is contained in:
AsherJingkongChen 2024-11-13 03:01:22 +08:00
commit f7d4c1ee1d
157 changed files with 11823 additions and 752 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -629,7 +629,7 @@ jobs:
run: taplo format --check --diff
- name: Check for typos
uses: crate-ci/typos@v1.26.8
uses: crate-ci/typos@v1.27.3
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.8
uses: JamesIves/github-pages-deploy-action@v4.6.9
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.8
uses: JamesIves/github-pages-deploy-action@v4.6.9
if: github.ref == 'refs/heads/trunk'
with:
token: ${{ secrets.WEB_DEPLOY }}

90
CHANGELOG.md
View File

@ -40,12 +40,69 @@ Bottom level categories:
## Unreleased
## Major changes
### The `diagnostic(…);` directive is now supported in WGSL
Naga now parses `diagnostic(…);` directives according to the WGSL spec. This allows users to control certain lints, similar to Rust's `allow`, `warn`, and `deny` attributes. For example, in standard WGSL (but, notably, not Naga yet—see <https://github.com/gfx-rs/wgpu/issues/4369>) this snippet would emit a uniformity error:
```wgsl
@group(0) @binding(0) var s : sampler;
@group(0) @binding(2) var tex : texture_2d<f32>;
@group(1) @binding(0) var<storage, read> ro_buffer : array<f32, 4>;
@fragment
fn main(@builtin(position) p : vec4f) -> @location(0) vec4f {
if ro_buffer[0] == 0 {
// Emits a derivative uniformity error during validation.
return textureSample(tex, s, vec2(0.,0.));
}
return vec4f(0.);
}
```
…but we can now silence it with the `off` severity level, like so:
```wgsl
// Disable the diagnosic with this…
diagnostic(off, derivative_uniformity);
@group(0) @binding(0) var s : sampler;
@group(0) @binding(2) var tex : texture_2d<f32>;
@group(1) @binding(0) var<storage, read> ro_buffer : array<f32, 4>;
@fragment
fn main(@builtin(position) p : vec4f) -> @location(0) vec4f {
if ro_buffer[0] == 0 {
// Look ma, no error!
return textureSample(tex, s, vec2(0.,0.));
}
return vec4f(0.);
}
```
There are some limitations to keep in mind with this new functionality:
- We do not yet support `diagnostic(…)` rules in attribute position (i.e., `@diagnostic(…) fn my_func { … }`). This is being tracked in <https://github.com/gfx-rs/wgpu/issues/5320>. We expect that rules in `fn` attribute position will be relaxed shortly (see <https://github.com/gfx-rs/wgpu/pull/6353>), but the prioritization for statement positions is unclear. If you are blocked by not being able to parse `diagnostic(…)` rules in statement positions, please let us know in that issue, so we can determine how to prioritize it!
- Standard WGSL specifies `error`, `warning`, `info`, and `off` severity levels. These are all technically usable now! A caveat, though: warning- and info-level are only emitted to `stderr` via the `log` façade, rather than being reported through a `Result::Err` in Naga or the `CompilationInfo` interface in `wgpu{,-core}`. This will require breaking changes in Naga to fix, and is being tracked by <https://github.com/gfx-rs/wgpu/issues/6458>.
- Not all lints can be controlled with `diagnostic(…)` rules. In fact, only the `derivative_uniformity` triggering rule exists in the WGSL standard. That said, Naga contributors are excited to see how this level of control unlocks a new ecosystem of configurable diagnostics.
- Finally, `diagnostic(…)` rules are not yet emitted in WGSL output. This means that `wgsl-in` → `wgsl-out` is currently a lossy process. We felt that it was important to unblock users who needed `diagnostic(…)` rules (i.e., <https://github.com/gfx-rs/wgpu/issues/3135>) before we took significant effort to fix this (tracked in <https://github.com/gfx-rs/wgpu/issues/6496>).
By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148](https://github.com/gfx-rs/wgpu/pull/6148).
### New Features
#### Naga
- Parse `diagnostic(…)` directives, but don't implement any triggering rules yet. By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456).
- Fix an issue where `naga` CLI would incorrectly skip the first positional argument when `--stdin-file-path` was specified. By @ErichDonGubler in [#6480](https://github.com/gfx-rs/wgpu/pull/6480).
- Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483).
- Implement `quantizeToF16()` for WGSL frontend, and WGSL, SPIR-V, HLSL, MSL, and GLSL backends. By @jamienicol in [#6519](https://github.com/gfx-rs/wgpu/pull/6519).
#### General
- Return submission index in `map_async` and `on_submitted_work_done` to track down completion of async callbacks. By @eliemichel in [#6360](https://github.com/gfx-rs/wgpu/pull/6360).
#### Vulkan
@ -57,10 +114,35 @@ Bottom level categories:
### Changes
#### Naga
- Show types of LHS and RHS in binary operation type mismatch errors. By @ErichDonGubler in [#6450](https://github.com/gfx-rs/wgpu/pull/6450).
#### General
- Make `Surface::as_hal` take an immutable reference to the surface. By @jerzywilczek in [#9999](https://github.com/gfx-rs/wgpu/pull/9999)
#### HAL
- Change the `DropCallback` API to use `FnOnce` instead of `FnMut`. By @jerzywilczek in [#6482](https://github.com/gfx-rs/wgpu/pull/6482)
### Bug Fixes
#### General
- Handle query set creation failure as an internal error that loses the `Device`, rather than panicking. By @ErichDonGubler in [#6505](https://github.com/gfx-rs/wgpu/pull/6505).
#### Naga
- Fix crash when a texture argument is missing. By @aedm in [#6486](https://github.com/gfx-rs/wgpu/pull/6486)
- Emit an error in constant evaluation, rather than crash, in certain cases where `vecN` constructors have less than N arguments. By @ErichDonGubler in [#6508](https://github.com/gfx-rs/wgpu/pull/6508).
#### General
- Ensure that `Features::TIMESTAMP_QUERY` is set when using timestamp writes in render and compute passes. By @ErichDonGubler in [#6497](https://github.com/gfx-rs/wgpu/pull/6497).
- Check for device mismatches when beginning render and compute passes. By @ErichDonGubler in [#6497](https://github.com/gfx-rs/wgpu/pull/6497).
- Lower `QUERY_SET_MAX_QUERIES` (and enforced limits) from 8192 to 4096 to match WebGPU spec. By @ErichDonGubler in [#6525](https://github.com/gfx-rs/wgpu/pull/6525).
## 23.0.0 (2024-10-25)
### Themes of this release
@ -75,7 +157,7 @@ This may not sound exciting, but let us convince you otherwise! All major web br
WGPU also benefits from standard, portable behavior in the same way as web browsers. Because of this behavior, it's generally fairly easy to port over usage of WebGPU in JavaScript to WGPU. It is also what lets WGPU go full circle: WGPU can be an implementation of WebGPU on native targets, but _also_ it can use _other implementations of WebGPU_ as a backend in JavaScript when compiled to WASM. Therefore, the same dynamic applies: if WGPU's own behavior were significantly different, then WGPU and end users would be _sad, sad humans_ as soon as they discover places where their nice apps are breaking, right?
The answer is: yes, we _do_ have sad, sad humans that really want their WGPU code to work _everywhere_. As Firefox and others use WGPU to implement WebGPU, the above example of Firefox diverging from standard is, unfortunately, today's reality. It _mostly_ behaves the same as a standards-compliant WebGPU, but it still doesn't in many important ways. Of particular note is Naga, its implementation of the WebGPU Shader Language. Shaders are pretty much a black-and-white point of failure in GPU programming; if they don't compile, then you can't use the rest of the API! And yet, it's extremely easy to run into this:
The answer is: yes, we _do_ have sad, sad humans that really want their WGPU code to work _everywhere_. As Firefox and others use WGPU to implement WebGPU, the above example of Firefox diverging from standard is, unfortunately, today's reality. It _mostly_ behaves the same as a standards-compliant WebGPU, but it still doesn't in many important ways. Of particular note is Naga, its implementation of the WebGPU Shader Language. Shaders are pretty much a black-and-white point of failure in GPU programming; if they don't compile, then you can't use the rest of the API! And yet, it's extremely easy to run into a case like that from <https://github.com/gfx-rs/wgpu/issues/4400>:
```wgsl
fn gimme_a_float() -> f32 {
@ -184,6 +266,10 @@ By @MarijnS95 in [#6006](https://github.com/gfx-rs/wgpu/pull/6006).
### New Features
#### Wgpu
- Added initial acceleration structure and ray query support into wgpu. By @expenses @daniel-keitel @Vecvec @JMS55 @atlv24 in [#6291](https://github.com/gfx-rs/wgpu/pull/6291)
#### Naga
- 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).

48
Cargo.lock generated
View File

@ -126,9 +126,9 @@ dependencies = [
[[package]]
name = "anstyle"
version = "1.0.9"
version = "1.0.10"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8365de52b16c035ff4fcafe0092ba9390540e3e352870ac09933bebcaa2c8c56"
checksum = "55cc3b69f167a1ef2e161439aa98aed94e6028e5f9a59be9a6ffb47aef1651f9"
[[package]]
name = "anstyle-parse"
@ -160,9 +160,9 @@ dependencies = [
[[package]]
name = "anyhow"
version = "1.0.91"
version = "1.0.92"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c042108f3ed77fd83760a5fd79b53be043192bb3b9dba91d8c574c0ada7850c8"
checksum = "74f37166d7d48a0284b99dd824694c26119c700b53bf0d1540cdb147dbdaaf13"
[[package]]
name = "arbitrary"
@ -434,9 +434,9 @@ checksum = "37b2a672a2cb129a2e41c10b1224bb368f9f37a2b16b612598138befd7b37eb5"
[[package]]
name = "cc"
version = "1.1.31"
version = "1.1.34"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c2e7962b54006dcfcc61cb72735f4d89bb97061dd6a7ed882ec6b8ee53714c6f"
checksum = "67b9470d453346108f93a59222a9a1a5724db32d0a4727b7ab7ace4b4d822dc9"
dependencies = [
"jobserver",
"libc",
@ -1195,9 +1195,9 @@ checksum = "9e5c1b78ca4aae1ac06c48a526a655760685149f0d465d21f37abfe57ce075c6"
[[package]]
name = "futures-lite"
version = "2.3.0"
version = "2.4.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "52527eb5074e35e9339c6b4e8d12600c7128b68fb25dcb9fa9dec18f7c25f3a5"
checksum = "3f1fa2f9765705486b33fd2acf1577f8ec449c2ba1f318ae5447697b7c08d210"
dependencies = [
"fastrand",
"futures-core",
@ -1305,6 +1305,9 @@ name = "glam"
version = "0.28.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "779ae4bf7e8421cf91c0b3b64e7e8b40b862fba4d393f59150042de7c4965a94"
dependencies = [
"bytemuck",
]
[[package]]
name = "glow"
@ -1700,7 +1703,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "4979f22fdb869068da03c9f7528f8297c6fd2606bc3a4affe42e6a823fdb8da4"
dependencies = [
"cfg-if",
"windows-targets 0.48.5",
"windows-targets 0.52.6",
]
[[package]]
@ -1834,9 +1837,9 @@ dependencies = [
[[package]]
name = "minicov"
version = "0.3.6"
version = "0.3.7"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "def6d99771d7c499c26ad4d40eb6645eafd3a1553b35fc26ea5a489a45e82d9a"
checksum = "f27fe9f1cc3c22e1687f9446c2083c4c5fc7f0bcf1c7a86bdbded14985895b4b"
dependencies = [
"cc",
"walkdir",
@ -2689,18 +2692,18 @@ checksum = "388a1df253eca08550bef6c72392cfe7c30914bf41df5269b68cbd6ff8f570a3"
[[package]]
name = "serde"
version = "1.0.213"
version = "1.0.214"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3ea7893ff5e2466df8d720bb615088341b295f849602c6956047f8f80f0e9bc1"
checksum = "f55c3193aca71c12ad7890f1785d2b73e1b9f63a0bbc353c08ef26fe03fc56b5"
dependencies = [
"serde_derive",
]
[[package]]
name = "serde_derive"
version = "1.0.213"
version = "1.0.214"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7e85ad2009c50b58e87caa8cd6dac16bdf511bbfb7af6c33df902396aa480fa5"
checksum = "de523f781f095e28fa605cdce0f8307e451cc0fd14e2eb4cd2e98a355b147766"
dependencies = [
"proc-macro2",
"quote",
@ -2927,9 +2930,9 @@ dependencies = [
[[package]]
name = "syn"
version = "2.0.85"
version = "2.0.87"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5023162dfcd14ef8f32034d8bcd4cc5ddc61ef7a247c024a33e24e1f24d21b56"
checksum = "25aa4ce346d03a6dcd68dd8b4010bcb74e54e62c90c573f394c46eae99aba32d"
dependencies = [
"proc-macro2",
"quote",
@ -2947,18 +2950,18 @@ dependencies = [
[[package]]
name = "thiserror"
version = "1.0.65"
version = "1.0.69"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5d11abd9594d9b38965ef50805c5e469ca9cc6f197f883f717e0269a3057b3d5"
checksum = "b6aaf5339b578ea85b50e080feb250a3e8ae8cfcdff9a461c9ec2904bc923f52"
dependencies = [
"thiserror-impl",
]
[[package]]
name = "thiserror-impl"
version = "1.0.65"
version = "1.0.69"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ae71770322cbd277e69d762a16c444af02aa0575ac0d174f0b9562d3b37f8602"
checksum = "4fee6c4efc90059e10f81e6d42c60a18f76588c3d74cb83a0b242a2b6c7504c1"
dependencies = [
"proc-macro2",
"quote",
@ -3746,6 +3749,7 @@ dependencies = [
"ctor",
"env_logger",
"futures-lite",
"glam",
"image",
"itertools",
"js-sys",
@ -3814,7 +3818,7 @@ version = "0.1.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cf221c93e13a30d793f7645a0e7762c55d169dbb0a49671918a2319d289b10bb"
dependencies = [
"windows-sys 0.48.0",
"windows-sys 0.59.0",
]
[[package]]

4
Cargo.toml
View File

@ -70,7 +70,7 @@ path = "./naga"
version = "23.0.0"
[workspace.dependencies]
anyhow = "1.0.91"
anyhow = "1.0.92"
argh = "0.1.5"
arrayvec = "0.7"
bincode = "1"
@ -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.65"
thiserror = "1.0.69"
wgpu = { version = "23.0.0", path = "./wgpu", default-features = false }
wgpu-core = { version = "23.0.0", path = "./wgpu-core" }
wgpu-macros = { version = "23.0.0", path = "./wgpu-macros" }

2
examples/Cargo.toml
View File

@ -34,7 +34,7 @@ cfg-if.workspace = true
encase = { workspace = true, features = ["glam"] }
flume.workspace = true
getrandom.workspace = true
glam.workspace = true
glam = { workspace = true, features = ["bytemuck"] }
ktx2.workspace = true
log.workspace = true
nanorand.workspace = true

84
examples/README.md
View File

@ -25,6 +25,9 @@ The rest of the examples are for demonstrating specific features that you can co
- `skybox` - Shows off too many concepts to list here. The name comes from game development where a "skybox" acts as a background for rendering, usually to add a sky texture for immersion, although they can also be used for backdrops to give the idea of a world beyond the game scene. This example does so much more than this, though, as it uses a car model loaded from a file and uses the user's mouse to rotate the car model in 3d. `skybox` also makes use of depth textures and similar app patterns to `uniform_values`.
- `shadow` - Likely by far the most complex example (certainly the largest in lines of code) of the official WGPU examples. `shadow` demonstrates basic scene rendering with the main attraction being lighting and shadows (as the name implies). It is recommended that any user looking into lighting be very familiar with the basic concepts of not only rendering with WGPU but also the primary mathematical ideas of computer graphics.
- `render_to_texture` - Renders to an image texture offscreen, demonstrating both off-screen rendering as well as how to add a sort of resolution-agnostic screenshot feature to an engine. This example either outputs an image file of your naming (pass command line arguments after specifying a `--` like `cargo run --bin wgpu-examples -- render_to_texture "test.png"`) or adds an `img` element containing the image to the page in WASM.
- `ray_cube_fragment` - Demonstrates using ray queries with a fragment shader.
- `ray_scene` - Demonstrates using ray queries and model loading
- `ray_shadows` - Demonstrates a simple use of ray queries - high quality shadows - uses a light set with push constants to raytrace through an untransformed scene and detect whether there is something obstructing the light.
#### Compute
@ -37,48 +40,51 @@ The rest of the examples are for demonstrating specific features that you can co
#### Combined
- `boids` - Demonstrates how to combine compute and render workflows by performing a [boid](https://en.wikipedia.org/wiki/Boids) simulation and rendering the boids to the screen as little triangles.
- `ray_cube_compute` - Demonstrates using ray queries with a compute shader.
- `ray_traced_triangle` - A simpler example demonstrating using ray queries with a compute shader
## Feature matrix
| Feature | boids | bunnymark | conservative_raster | cube | hello_synchronization | hello_workgroups | mipmap | msaa_line | render_to_texture | repeated_compute | shadow | skybox | stencil_triangles | storage_texture | texture_arrays | uniform_values | water |
| ---------------------------- | ------ | --------- | ------------------- | ------ | --------------------- | ---------------- | ------ | --------- | ----------------- | ---------------- | ------ | ------ | ----------------- | --------------- | -------------- | -------------- | ------ |
| vertex attributes | :star: | | | :star: | | | | :star: | | | :star: | :star: | | | :star: | | :star: |
| instancing | :star: | | | | | | | | | | | | | | | | |
| lines and points | | | :star: | | | | | :star: | | | | | | | | | |
| dynamic buffer offsets | | :star: | | | | | | | | | :star: | | | | | | |
| implicit layout | | | | | | | :star: | | | | | | | | | | |
| sampled color textures | :star: | :star: | :star: | :star: | | | :star: | | | | | :star: | | | :star: | | :star: |
| storage textures | :star: | | | | | | | | | | | | | :star: | | | |
| comparison samplers | | | | | | | | | | | :star: | | | | | | |
| subresource views | | | | | | | :star: | | | | :star: | | | | | | |
| cubemaps | | | | | | | | | | | | :star: | | | | | |
| multisampling | | | | | | | | :star: | | | | | | | | | |
| off-screen rendering | | | :star: | | | | | | :star: | | :star: | | | | | | :star: |
| stencil testing | | | | | | | | | | | | | :star: | | | | |
| depth testing | | | | | | | | | | | :star: | :star: | | | | | :star: |
| depth biasing | | | | | | | | | | | :star: | | | | | | |
| read-only depth | | | | | | | | | | | | | | | | | :star: |
| blending | | :star: | | :star: | | | | | | | | | | | | | :star: |
| render bundles | | | | | | | | :star: | | | | | | | | | :star: |
| uniform buffers | | | | | | | | | | | | | | | | :star: | |
| compute passes | :star: | | | | :star: | :star: | | | | :star: | | | | :star: | | | |
| buffer mapping | | | | | :star: | :star: | | | | :star: | | | | :star: | | | |
| error scopes | | | | :star: | | | | | | | | | | | | | |
| compute workgroups | | | | | :star: | :star: | | | | | | | | | | | |
| compute synchronization | | | | | :star: | | | | | | | | | | | | |
| _optional extensions_ | | | | | | | | | | | | | | | :star: | | |
| - SPIR-V shaders | | | | | | | | | | | | | | | | | |
| - binding array | | | | | | | | | | | | | | | :star: | | |
| - push constants | | | | | | | | | | | | | | | | | |
| - depth clamping | | | | | | | | | | | :star: | | | | | | |
| - compressed textures | | | | | | | | | | | | :star: | | | | | |
| - polygon mode | | | | :star: | | | | | | | | | | | | | |
| - queries | | | | | | | :star: | | | | | | | | | | |
| - conservative rasterization | | | :star: | | | | | | | | | | | | | | |
| _integrations_ | | | | | | | | | | | | | | | | | |
| - staging belt | | | | | | | | | | | | :star: | | | | | |
| - typed arena | | | | | | | | | | | | | | | | | |
| - obj loading | | | | | | | | | | | | :star: | | | | | |
| Feature | boids | bunnymark | conservative_raster | cube | hello_synchronization | hello_workgroups | mipmap | msaa_line | render_to_texture | repeated_compute | shadow | skybox | stencil_triangles | storage_texture | texture_arrays | uniform_values | water | ray_cube_compute | ray_cube_fragment | ray_scene | ray_shadows | ray_traced_triangle |
|------------------------------| ------ | --------- | ------------------- | ------ | --------------------- | ---------------- | ------ | --------- | ----------------- | ---------------- | ------ | ------ | ----------------- | --------------- | -------------- | -------------- | ------ |------------------|-------------------|-----------|-------------|---------------------|
| vertex attributes | :star: | | | :star: | | | | :star: | | | :star: | :star: | | | :star: | | :star: | | | | | |
| instancing | :star: | | | | | | | | | | | | | | | | | | | | | |
| lines and points | | | :star: | | | | | :star: | | | | | | | | | | | | | | |
| dynamic buffer offsets | | :star: | | | | | | | | | :star: | | | | | | | | | | | |
| implicit layout | | | | | | | :star: | | | | | | | | | | | | | | | |
| sampled color textures | :star: | :star: | :star: | :star: | | | :star: | | | | | :star: | | | :star: | | :star: | | | | | |
| storage textures | :star: | | | | | | | | | | | | | :star: | | | | :star: | | | | :star: |
| comparison samplers | | | | | | | | | | | :star: | | | | | | | | | | | |
| subresource views | | | | | | | :star: | | | | :star: | | | | | | | | | | | |
| cubemaps | | | | | | | | | | | | :star: | | | | | | | | | | |
| multisampling | | | | | | | | :star: | | | | | | | | | | | | | | |
| off-screen rendering | | | :star: | | | | | | :star: | | :star: | | | | | | :star: | | | | | |
| stencil testing | | | | | | | | | | | | | :star: | | | | | | | | | |
| depth testing | | | | | | | | | | | :star: | :star: | | | | | :star: | | | | | |
| depth biasing | | | | | | | | | | | :star: | | | | | | | | | | | |
| read-only depth | | | | | | | | | | | | | | | | | :star: | | | | | |
| blending | | :star: | | :star: | | | | | | | | | | | | | :star: | | | | | |
| render bundles | | | | | | | | :star: | | | | | | | | | :star: | | | | | |
| uniform buffers | | | | | | | | | | | | | | | | :star: | | | | | | |
| compute passes | :star: | | | | :star: | :star: | | | | :star: | | | | :star: | | | | | | | | |
| buffer mapping | | | | | :star: | :star: | | | | :star: | | | | :star: | | | | | | | | |
| error scopes | | | | :star: | | | | | | | | | | | | | | | | | | |
| compute workgroups | | | | | :star: | :star: | | | | | | | | | | | | | | | | |
| compute synchronization | | | | | :star: | | | | | | | | | | | | | | | | | |
| _optional extensions_ | | | | | | | | | | | | | | | :star: | | | | | | | |
| - SPIR-V shaders | | | | | | | | | | | | | | | | | | | | | | |
| - binding array | | | | | | | | | | | | | | | :star: | | | | | | | |
| - push constants | | | | | | | | | | | | | | | | | | | | | :star: | |
| - depth clamping | | | | | | | | | | | :star: | | | | | | | | | | | |
| - compressed textures | | | | | | | | | | | | :star: | | | | | | | | | | |
| - polygon mode | | | | :star: | | | | | | | | | | | | | | | | | | |
| - queries | | | | | | | :star: | | | | | | | | | | | | | | | |
| - conservative rasterization | | | :star: | | | | | | | | | | | | | | | | | | | |
| - ray queries | | | | | | | | | | | | | | | | | | :star: | :star: | :star: | :star: | :star: |
| _integrations_ | | | | | | | | | | | | | | | | | | | | | | |
| - staging belt | | | | | | | | | | | | :star: | | | | | | | | | | |
| - typed arena | | | | | | | | | | | | | | | | | | | | | | |
| - obj loading | | | | | | | | | | | | :star: | | | | | | | | :star: | | |
## Running on the Web

5
examples/src/lib.rs
View File

@ -15,6 +15,11 @@ pub mod hello_windows;
pub mod hello_workgroups;
pub mod mipmap;
pub mod msaa_line;
pub mod ray_cube_compute;
pub mod ray_cube_fragment;
pub mod ray_scene;
pub mod ray_shadows;
pub mod ray_traced_triangle;
pub mod render_to_texture;
pub mod repeated_compute;
pub mod shadow;

30
examples/src/main.rs
View File

@ -146,6 +146,36 @@ const EXAMPLES: &[ExampleDesc] = &[
webgl: false, // No RODS
webgpu: true,
},
ExampleDesc {
name: "ray_cube_compute",
function: wgpu_examples::ray_cube_compute::main,
webgl: false, // No Ray-tracing extensions
webgpu: false, // No Ray-tracing extensions (yet)
},
ExampleDesc {
name: "ray_cube_fragment",
function: wgpu_examples::ray_cube_fragment::main,
webgl: false, // No Ray-tracing extensions
webgpu: false, // No Ray-tracing extensions (yet)
},
ExampleDesc {
name: "ray_scene",
function: wgpu_examples::ray_scene::main,
webgl: false, // No Ray-tracing extensions
webgpu: false, // No Ray-tracing extensions (yet)
},
ExampleDesc {
name: "ray_shadows",
function: wgpu_examples::ray_shadows::main,
webgl: false, // No Ray-tracing extensions
webgpu: false, // No Ray-tracing extensions (yet)
},
ExampleDesc {
name: "ray_traced_triangle",
function: wgpu_examples::ray_traced_triangle::main,
webgl: false,
webgpu: false,
},
];
fn get_example_name() -> Option<String> {

14
examples/src/ray_cube_compute/README.md Normal file
View File

@ -0,0 +1,14 @@
# ray-cube
This example renders a ray traced cube with hardware acceleration.
A separate compute shader is used to perform the ray queries.
## To Run
```
cargo run --bin wgpu-examples ray_cube_compute
```
## Screenshots
![Cube example](screenshot.png)

52
examples/src/ray_cube_compute/blit.wgsl Normal file
View File

@ -0,0 +1,52 @@
struct VertexOutput {
@builtin(position) position: vec4<f32>,
@location(0) tex_coords: vec2<f32>,
};
// meant to be called with 3 vertex indices: 0, 1, 2
// draws one large triangle over the clip space like this:
// (the asterisks represent the clip space bounds)
//-1,1 1,1
// ---------------------------------
// | * .
// | * .
// | * .
// | * .
// | * .
// | * .
// |***************
// | . 1,-1
// | .
// | .
// | .
// | .
// |.
@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
var result: VertexOutput;
let x = i32(vertex_index) / 2;
let y = i32(vertex_index) & 1;
let tc = vec2<f32>(
f32(x) * 2.0,
f32(y) * 2.0
);
result.position = vec4<f32>(
tc.x * 2.0 - 1.0,
1.0 - tc.y * 2.0,
0.0, 1.0
);
result.tex_coords = tc;
return result;
}
@group(0)
@binding(0)
var r_color: texture_2d<f32>;
@group(0)
@binding(1)
var r_sampler: sampler;
@fragment
fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
return textureSample(r_color, r_sampler, vertex.tex_coords);
}

500
examples/src/ray_cube_compute/mod.rs Normal file
View File

@ -0,0 +1,500 @@
use std::{borrow::Cow, future::Future, iter, mem, pin::Pin, task, time::Instant};
use bytemuck::{Pod, Zeroable};
use glam::{Affine3A, Mat4, Quat, Vec3};
use wgpu::util::DeviceExt;
use wgpu::StoreOp;
// from cube
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Vertex {
_pos: [f32; 4],
_tex_coord: [f32; 2],
}
fn vertex(pos: [i8; 3], tc: [i8; 2]) -> Vertex {
Vertex {
_pos: [pos[0] as f32, pos[1] as f32, pos[2] as f32, 1.0],
_tex_coord: [tc[0] as f32, tc[1] as f32],
}
}
fn create_vertices() -> (Vec<Vertex>, Vec<u16>) {
let vertex_data = [
// top (0, 0, 1)
vertex([-1, -1, 1], [0, 0]),
vertex([1, -1, 1], [1, 0]),
vertex([1, 1, 1], [1, 1]),
vertex([-1, 1, 1], [0, 1]),
// bottom (0, 0, -1)
vertex([-1, 1, -1], [1, 0]),
vertex([1, 1, -1], [0, 0]),
vertex([1, -1, -1], [0, 1]),
vertex([-1, -1, -1], [1, 1]),
// right (1, 0, 0)
vertex([1, -1, -1], [0, 0]),
vertex([1, 1, -1], [1, 0]),
vertex([1, 1, 1], [1, 1]),
vertex([1, -1, 1], [0, 1]),
// left (-1, 0, 0)
vertex([-1, -1, 1], [1, 0]),
vertex([-1, 1, 1], [0, 0]),
vertex([-1, 1, -1], [0, 1]),
vertex([-1, -1, -1], [1, 1]),
// front (0, 1, 0)
vertex([1, 1, -1], [1, 0]),
vertex([-1, 1, -1], [0, 0]),
vertex([-1, 1, 1], [0, 1]),
vertex([1, 1, 1], [1, 1]),
// back (0, -1, 0)
vertex([1, -1, 1], [0, 0]),
vertex([-1, -1, 1], [1, 0]),
vertex([-1, -1, -1], [1, 1]),
vertex([1, -1, -1], [0, 1]),
];
let index_data: &[u16] = &[
0, 1, 2, 2, 3, 0, // top
4, 5, 6, 6, 7, 4, // bottom
8, 9, 10, 10, 11, 8, // right
12, 13, 14, 14, 15, 12, // left
16, 17, 18, 18, 19, 16, // front
20, 21, 22, 22, 23, 20, // back
];
(vertex_data.to_vec(), index_data.to_vec())
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Uniforms {
view_inverse: Mat4,
proj_inverse: Mat4,
}
#[inline]
fn affine_to_rows(mat: &Affine3A) -> [f32; 12] {
let row_0 = mat.matrix3.row(0);
let row_1 = mat.matrix3.row(1);
let row_2 = mat.matrix3.row(2);
let translation = mat.translation;
[
row_0.x,
row_0.y,
row_0.z,
translation.x,
row_1.x,
row_1.y,
row_1.z,
translation.y,
row_2.x,
row_2.y,
row_2.z,
translation.z,
]
}
/// A wrapper for `pop_error_scope` futures that panics if an error occurs.
///
/// Given a future `inner` of an `Option<E>` for some error type `E`,
/// wait for the future to be ready, and panic if its value is `Some`.
///
/// This can be done simpler with `FutureExt`, but we don't want to add
/// a dependency just for this small case.
struct ErrorFuture<F> {
inner: F,
}
impl<F: Future<Output = Option<wgpu::Error>>> Future for ErrorFuture<F> {
type Output = ();
fn poll(self: Pin<&mut Self>, cx: &mut task::Context<'_>) -> task::Poll<()> {
let inner = unsafe { self.map_unchecked_mut(|me| &mut me.inner) };
inner.poll(cx).map(|error| {
if let Some(e) = error {
panic!("Rendering {}", e);
}
})
}
}
#[allow(dead_code)]
struct Example {
rt_target: wgpu::Texture,
rt_view: wgpu::TextureView,
sampler: wgpu::Sampler,
uniform_buf: wgpu::Buffer,
vertex_buf: wgpu::Buffer,
index_buf: wgpu::Buffer,
tlas_package: wgpu::TlasPackage,
compute_pipeline: wgpu::ComputePipeline,
compute_bind_group: wgpu::BindGroup,
blit_pipeline: wgpu::RenderPipeline,
blit_bind_group: wgpu::BindGroup,
start_inst: Instant,
}
impl crate::framework::Example for Example {
fn required_features() -> wgpu::Features {
wgpu::Features::TEXTURE_BINDING_ARRAY
| wgpu::Features::STORAGE_RESOURCE_BINDING_ARRAY
| wgpu::Features::VERTEX_WRITABLE_STORAGE
| wgpu::Features::EXPERIMENTAL_RAY_QUERY
| wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
}
fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities {
wgpu::DownlevelCapabilities::default()
}
fn required_limits() -> wgpu::Limits {
wgpu::Limits::default()
}
fn init(
config: &wgpu::SurfaceConfiguration,
_adapter: &wgpu::Adapter,
device: &wgpu::Device,
queue: &wgpu::Queue,
) -> Self {
let side_count = 8;
let rt_target = device.create_texture(&wgpu::TextureDescriptor {
label: Some("rt_target"),
size: wgpu::Extent3d {
width: config.width,
height: config.height,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba8Unorm,
usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING,
view_formats: &[wgpu::TextureFormat::Rgba8Unorm],
});
let rt_view = rt_target.create_view(&wgpu::TextureViewDescriptor {
label: None,
format: Some(wgpu::TextureFormat::Rgba8Unorm),
dimension: Some(wgpu::TextureViewDimension::D2),
aspect: wgpu::TextureAspect::All,
base_mip_level: 0,
mip_level_count: None,
base_array_layer: 0,
array_layer_count: None,
});
let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
label: Some("rt_sampler"),
address_mode_u: wgpu::AddressMode::ClampToEdge,
address_mode_v: wgpu::AddressMode::ClampToEdge,
address_mode_w: wgpu::AddressMode::ClampToEdge,
mag_filter: wgpu::FilterMode::Linear,
min_filter: wgpu::FilterMode::Linear,
mipmap_filter: wgpu::FilterMode::Nearest,
..Default::default()
});
let uniforms = {
let view = Mat4::look_at_rh(Vec3::new(0.0, 0.0, 2.5), Vec3::ZERO, Vec3::Y);
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.001,
1000.0,
);
Uniforms {
view_inverse: view.inverse(),
proj_inverse: proj.inverse(),
}
};
let uniform_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Uniform Buffer"),
contents: bytemuck::cast_slice(&[uniforms]),
usage: wgpu::BufferUsages::UNIFORM,
});
let (vertex_data, index_data) = create_vertices();
let vertex_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Vertex Buffer"),
contents: bytemuck::cast_slice(&vertex_data),
usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::BLAS_INPUT,
});
let index_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Index Buffer"),
contents: bytemuck::cast_slice(&index_data),
usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::BLAS_INPUT,
});
let blas_geo_size_desc = wgpu::BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
vertex_count: vertex_data.len() as u32,
index_format: Some(wgpu::IndexFormat::Uint16),
index_count: Some(index_data.len() as u32),
flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
};
let blas = device.create_blas(
&wgpu::CreateBlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
},
wgpu::BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_geo_size_desc.clone()],
},
);
let tlas = device.create_tlas(&wgpu::CreateTlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
max_instances: side_count * side_count,
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("rt_computer"),
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let blit_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("blit"),
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("blit.wgsl"))),
});
let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("rt"),
layout: None,
module: &shader,
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
let compute_bind_group_layout = compute_pipeline.get_bind_group_layout(0);
let compute_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &compute_bind_group_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(&rt_view),
},
wgpu::BindGroupEntry {
binding: 1,
resource: uniform_buf.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 2,
resource: wgpu::BindingResource::AccelerationStructure(&tlas),
},
],
});
let blit_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("blit"),
layout: None,
vertex: wgpu::VertexState {
module: &blit_shader,
entry_point: Some("vs_main"),
compilation_options: Default::default(),
buffers: &[],
},
fragment: Some(wgpu::FragmentState {
module: &blit_shader,
entry_point: Some("fs_main"),
compilation_options: Default::default(),
targets: &[Some(config.format.into())],
}),
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleList,
..Default::default()
},
depth_stencil: None,
multisample: wgpu::MultisampleState::default(),
multiview: None,
cache: None,
});
let blit_bind_group_layout = blit_pipeline.get_bind_group_layout(0);
let blit_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &blit_bind_group_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(&rt_view),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::Sampler(&sampler),
},
],
});
let mut tlas_package = wgpu::TlasPackage::new(tlas);
let dist = 3.0;
for x in 0..side_count {
for y in 0..side_count {
tlas_package[(x + y * side_count) as usize] = Some(wgpu::TlasInstance::new(
&blas,
affine_to_rows(&Affine3A::from_rotation_translation(
Quat::from_rotation_y(45.9_f32.to_radians()),
Vec3 {
x: x as f32 * dist,
y: y as f32 * dist,
z: -30.0,
},
)),
0,
0xff,
));
}
}
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(
iter::once(&wgpu::BlasBuildEntry {
blas: &blas,
geometry: wgpu::BlasGeometries::TriangleGeometries(vec![
wgpu::BlasTriangleGeometry {
size: &blas_geo_size_desc,
vertex_buffer: &vertex_buf,
first_vertex: 0,
vertex_stride: mem::size_of::<Vertex>() as u64,
index_buffer: Some(&index_buf),
index_buffer_offset: Some(0),
transform_buffer: None,
transform_buffer_offset: None,
},
]),
}),
iter::once(&tlas_package),
);
queue.submit(Some(encoder.finish()));
let start_inst = Instant::now();
Example {
rt_target,
rt_view,
sampler,
uniform_buf,
vertex_buf,
index_buf,
tlas_package,
compute_pipeline,
compute_bind_group,
blit_pipeline,
blit_bind_group,
start_inst,
}
}
fn update(&mut self, _event: winit::event::WindowEvent) {
//empty
}
fn resize(
&mut self,
_config: &wgpu::SurfaceConfiguration,
_device: &wgpu::Device,
_queue: &wgpu::Queue,
) {
}
fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
device.push_error_scope(wgpu::ErrorFilter::Validation);
let anim_time = self.start_inst.elapsed().as_secs_f64() as f32;
self.tlas_package[0].as_mut().unwrap().transform =
affine_to_rows(&Affine3A::from_rotation_translation(
Quat::from_euler(
glam::EulerRot::XYZ,
anim_time * 0.342,
anim_time * 0.254,
anim_time * 0.832,
),
Vec3 {
x: 0.0,
y: 0.0,
z: -6.0,
},
));
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(iter::empty(), iter::once(&self.tlas_package));
{
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
cpass.set_pipeline(&self.compute_pipeline);
cpass.set_bind_group(0, Some(&self.compute_bind_group), &[]);
cpass.dispatch_workgroups(self.rt_target.width() / 8, self.rt_target.height() / 8, 1);
}
{
let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view,
resolve_target: None,
ops: wgpu::Operations {
load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
store: StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
rpass.set_pipeline(&self.blit_pipeline);
rpass.set_bind_group(0, Some(&self.blit_bind_group), &[]);
rpass.draw(0..3, 0..1);
}
queue.submit(Some(encoder.finish()));
}
}
pub fn main() {
crate::framework::run::<Example>("ray-cube");
}
#[cfg(test)]
#[wgpu_test::gpu_test]
static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
name: "ray_cube_compute",
image_path: "/examples/src/ray_cube_compute/screenshot.png",
width: 1024,
height: 768,
optional_features: wgpu::Features::default(),
base_test_parameters: wgpu_test::TestParameters {
required_features: <Example as crate::framework::Example>::required_features(),
required_limits: <Example as crate::framework::Example>::required_limits(),
force_fxc: false,
skips: vec![],
failures: Vec::new(),
required_downlevel_caps:
<Example as crate::framework::Example>::required_downlevel_capabilities(),
},
comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
_phantom: std::marker::PhantomData::<Example>,
};

BIN
examples/src/ray_cube_compute/screenshot.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 479 KiB

82
examples/src/ray_cube_compute/shader.wgsl Normal file
View File

@ -0,0 +1,82 @@
/*
The contents of the RayQuery struct are roughly as follows
let RAY_FLAG_NONE = 0x00u;
let RAY_FLAG_OPAQUE = 0x01u;
let RAY_FLAG_NO_OPAQUE = 0x02u;
let RAY_FLAG_TERMINATE_ON_FIRST_HIT = 0x04u;
let RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08u;
let RAY_FLAG_CULL_BACK_FACING = 0x10u;
let RAY_FLAG_CULL_FRONT_FACING = 0x20u;
let RAY_FLAG_CULL_OPAQUE = 0x40u;
let RAY_FLAG_CULL_NO_OPAQUE = 0x80u;
let RAY_FLAG_SKIP_TRIANGLES = 0x100u;
let RAY_FLAG_SKIP_AABBS = 0x200u;
let RAY_QUERY_INTERSECTION_NONE = 0u;
let RAY_QUERY_INTERSECTION_TRIANGLE = 1u;
let RAY_QUERY_INTERSECTION_GENERATED = 2u;
let RAY_QUERY_INTERSECTION_AABB = 4u;
struct RayDesc {
flags: u32,
cull_mask: u32,
t_min: f32,
t_max: f32,
origin: vec3<f32>,
dir: vec3<f32>,
}
struct RayIntersection {
kind: u32,
t: f32,
instance_custom_index: u32,
instance_id: u32,
sbt_record_offset: u32,
geometry_index: u32,
primitive_index: u32,
barycentrics: vec2<f32>,
front_face: bool,
object_to_world: mat4x3<f32>,
world_to_object: mat4x3<f32>,
}
*/
struct Uniforms {
view_inv: mat4x4<f32>,
proj_inv: mat4x4<f32>,
};
@group(0) @binding(0)
var output: texture_storage_2d<rgba8unorm, write>;
@group(0) @binding(1)
var<uniform> uniforms: Uniforms;
@group(0) @binding(2)
var acc_struct: acceleration_structure;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
let target_size = textureDimensions(output);
var color = vec4<f32>(vec2<f32>(global_id.xy) / vec2<f32>(target_size), 0.0, 1.0);
let pixel_center = vec2<f32>(global_id.xy) + vec2<f32>(0.5);
let in_uv = pixel_center/vec2<f32>(target_size.xy);
let d = in_uv * 2.0 - 1.0;
let origin = (uniforms.view_inv * vec4<f32>(0.0,0.0,0.0,1.0)).xyz;
let temp = uniforms.proj_inv * vec4<f32>(d.x, d.y, 1.0, 1.0);
let direction = (uniforms.view_inv * vec4<f32>(normalize(temp.xyz), 0.0)).xyz;
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
color = vec4<f32>(intersection.barycentrics, 1.0 - intersection.barycentrics.x - intersection.barycentrics.y, 1.0);
}
textureStore(output, global_id.xy, color);
}

13
examples/src/ray_cube_fragment/README.md Normal file
View File

@ -0,0 +1,13 @@
# ray-cube
This example renders a ray traced cube with hardware acceleration.
## To Run
```
cargo run --bin wgpu-examples ray_cube_fragment
```
## Screenshots
![Cube example](screenshot.png)

391
examples/src/ray_cube_fragment/mod.rs Normal file
View File

@ -0,0 +1,391 @@
use bytemuck::{Pod, Zeroable};
use glam::{Mat4, Quat, Vec3};
use std::ops::IndexMut;
use std::{borrow::Cow, future::Future, iter, mem, pin::Pin, task, time::Instant};
use wgpu::util::DeviceExt;
// from cube
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Vertex {
_pos: [f32; 4],
_tex_coord: [f32; 2],
}
fn vertex(pos: [i8; 3], tc: [i8; 2]) -> Vertex {
Vertex {
_pos: [pos[0] as f32, pos[1] as f32, pos[2] as f32, 1.0],
_tex_coord: [tc[0] as f32, tc[1] as f32],
}
}
fn create_vertices() -> (Vec<Vertex>, Vec<u16>) {
let vertex_data = [
// top (0, 0, 1)
vertex([-1, -1, 1], [0, 0]),
vertex([1, -1, 1], [1, 0]),
vertex([1, 1, 1], [1, 1]),
vertex([-1, 1, 1], [0, 1]),
// bottom (0, 0, -1)
vertex([-1, 1, -1], [1, 0]),
vertex([1, 1, -1], [0, 0]),
vertex([1, -1, -1], [0, 1]),
vertex([-1, -1, -1], [1, 1]),
// right (1, 0, 0)
vertex([1, -1, -1], [0, 0]),
vertex([1, 1, -1], [1, 0]),
vertex([1, 1, 1], [1, 1]),
vertex([1, -1, 1], [0, 1]),
// left (-1, 0, 0)
vertex([-1, -1, 1], [1, 0]),
vertex([-1, 1, 1], [0, 0]),
vertex([-1, 1, -1], [0, 1]),
vertex([-1, -1, -1], [1, 1]),
// front (0, 1, 0)
vertex([1, 1, -1], [1, 0]),
vertex([-1, 1, -1], [0, 0]),
vertex([-1, 1, 1], [0, 1]),
vertex([1, 1, 1], [1, 1]),
// back (0, -1, 0)
vertex([1, -1, 1], [0, 0]),
vertex([-1, -1, 1], [1, 0]),
vertex([-1, -1, -1], [1, 1]),
vertex([1, -1, -1], [0, 1]),
];
let index_data: &[u16] = &[
0, 1, 2, 2, 3, 0, // top
4, 5, 6, 6, 7, 4, // bottom
8, 9, 10, 10, 11, 8, // right
12, 13, 14, 14, 15, 12, // left
16, 17, 18, 18, 19, 16, // front
20, 21, 22, 22, 23, 20, // back
];
(vertex_data.to_vec(), index_data.to_vec())
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Uniforms {
view_inverse: Mat4,
proj_inverse: Mat4,
}
/// A wrapper for `pop_error_scope` futures that panics if an error occurs.
///
/// Given a future `inner` of an `Option<E>` for some error type `E`,
/// wait for the future to be ready, and panic if its value is `Some`.
///
/// This can be done simpler with `FutureExt`, but we don't want to add
/// a dependency just for this small case.
struct ErrorFuture<F> {
inner: F,
}
impl<F: Future<Output = Option<wgpu::Error>>> Future for ErrorFuture<F> {
type Output = ();
fn poll(self: Pin<&mut Self>, cx: &mut task::Context<'_>) -> task::Poll<()> {
let inner = unsafe { self.map_unchecked_mut(|me| &mut me.inner) };
inner.poll(cx).map(|error| {
if let Some(e) = error {
panic!("Rendering {}", e);
}
})
}
}
struct Example {
uniforms: Uniforms,
uniform_buf: wgpu::Buffer,
blas: wgpu::Blas,
tlas_package: wgpu::TlasPackage,
pipeline: wgpu::RenderPipeline,
bind_group: wgpu::BindGroup,
start_inst: Instant,
}
impl crate::framework::Example for Example {
fn required_features() -> wgpu::Features {
wgpu::Features::EXPERIMENTAL_RAY_QUERY
| wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
}
fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities {
wgpu::DownlevelCapabilities::default()
}
fn required_limits() -> wgpu::Limits {
wgpu::Limits::default()
}
fn init(
config: &wgpu::SurfaceConfiguration,
_adapter: &wgpu::Adapter,
device: &wgpu::Device,
queue: &wgpu::Queue,
) -> Self {
let side_count = 8;
let uniforms = {
let view = Mat4::look_at_rh(Vec3::new(0.0, 0.0, 2.5), Vec3::ZERO, Vec3::Y);
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.001,
1000.0,
);
Uniforms {
view_inverse: view.inverse(),
proj_inverse: proj.inverse(),
}
};
let uniform_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Uniform Buffer"),
contents: bytemuck::cast_slice(&[uniforms]),
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
});
let (vertex_data, index_data) = create_vertices();
let vertex_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Vertex Buffer"),
contents: bytemuck::cast_slice(&vertex_data),
usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::BLAS_INPUT,
});
let index_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Index Buffer"),
contents: bytemuck::cast_slice(&index_data),
usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::BLAS_INPUT,
});
let blas_geo_size_desc = wgpu::BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
vertex_count: vertex_data.len() as u32,
index_format: Some(wgpu::IndexFormat::Uint16),
index_count: Some(index_data.len() as u32),
flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
};
let blas = device.create_blas(
&wgpu::CreateBlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
},
wgpu::BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_geo_size_desc.clone()],
},
);
let tlas = device.create_tlas(&wgpu::CreateTlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
max_instances: side_count * side_count,
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: None,
layout: None,
vertex: wgpu::VertexState {
module: &shader,
entry_point: Some("vs_main"),
compilation_options: Default::default(),
buffers: &[],
},
fragment: Some(wgpu::FragmentState {
module: &shader,
entry_point: Some("fs_main"),
compilation_options: Default::default(),
targets: &[Some(config.format.into())],
}),
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleList,
..Default::default()
},
depth_stencil: None,
multisample: wgpu::MultisampleState::default(),
multiview: None,
cache: None,
});
let bind_group_layout = pipeline.get_bind_group_layout(0);
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &bind_group_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: uniform_buf.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::AccelerationStructure(&tlas),
},
],
});
let tlas_package = wgpu::TlasPackage::new(tlas);
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(
iter::once(&wgpu::BlasBuildEntry {
blas: &blas,
geometry: wgpu::BlasGeometries::TriangleGeometries(vec![
wgpu::BlasTriangleGeometry {
size: &blas_geo_size_desc,
vertex_buffer: &vertex_buf,
first_vertex: 0,
vertex_stride: mem::size_of::<Vertex>() as u64,
index_buffer: Some(&index_buf),
index_buffer_offset: Some(0),
transform_buffer: None,
transform_buffer_offset: None,
},
]),
}),
// iter::empty(),
iter::once(&tlas_package),
);
queue.submit(Some(encoder.finish()));
let start_inst = Instant::now();
Example {
uniforms,
uniform_buf,
blas,
tlas_package,
pipeline,
bind_group,
start_inst,
}
}
fn update(&mut self, _event: winit::event::WindowEvent) {}
fn resize(
&mut self,
config: &wgpu::SurfaceConfiguration,
_device: &wgpu::Device,
queue: &wgpu::Queue,
) {
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.001,
1000.0,
);
self.uniforms.proj_inverse = proj.inverse();
queue.write_buffer(&self.uniform_buf, 0, bytemuck::cast_slice(&[self.uniforms]));
}
fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
device.push_error_scope(wgpu::ErrorFilter::Validation);
// scene update
{
let dist = 12.0;
let side_count = 8;
let anim_time = self.start_inst.elapsed().as_secs_f64() as f32;
for x in 0..side_count {
for y in 0..side_count {
let instance = self.tlas_package.index_mut((x + y * side_count) as usize);
let x = x as f32 / (side_count - 1) as f32;
let y = y as f32 / (side_count - 1) as f32;
let x = x * 2.0 - 1.0;
let y = y * 2.0 - 1.0;
let transform = Mat4::from_rotation_translation(
Quat::from_euler(
glam::EulerRot::XYZ,
anim_time * 0.5 * 0.342,
anim_time * 0.5 * 0.254,
anim_time * 0.5 * 0.832,
),
Vec3 {
x: x * dist,
y: y * dist,
z: -24.0,
},
);
let transform = transform.transpose().to_cols_array()[..12]
.try_into()
.unwrap();
*instance = Some(wgpu::TlasInstance::new(&self.blas, transform, 0, 0xff));
}
}
}
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(iter::empty(), iter::once(&self.tlas_package));
{
let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view,
resolve_target: None,
ops: wgpu::Operations {
load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
store: wgpu::StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
rpass.set_pipeline(&self.pipeline);
rpass.set_bind_group(0, Some(&self.bind_group), &[]);
rpass.draw(0..3, 0..1);
}
queue.submit(Some(encoder.finish()));
}
}
pub fn main() {
crate::framework::run::<Example>("ray-cube");
}
#[cfg(test)]
#[wgpu_test::gpu_test]
static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
name: "ray_cube_fragment",
image_path: "/examples/src/ray_cube_fragment/screenshot.png",
width: 1024,
height: 768,
optional_features: wgpu::Features::default(),
base_test_parameters: wgpu_test::TestParameters {
required_features: <Example as crate::framework::Example>::required_features(),
required_limits: <Example as crate::framework::Example>::required_limits(),
force_fxc: false,
skips: vec![],
failures: Vec::new(),
required_downlevel_caps:
<Example as crate::framework::Example>::required_downlevel_capabilities(),
},
comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
_phantom: std::marker::PhantomData::<Example>,
};

BIN
examples/src/ray_cube_fragment/screenshot.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 584 KiB

74
examples/src/ray_cube_fragment/shader.wgsl Normal file
View File

@ -0,0 +1,74 @@
struct VertexOutput {
@builtin(position) position: vec4<f32>,
@location(0) tex_coords: vec2<f32>,
};
// meant to be called with 3 vertex indices: 0, 1, 2
// draws one large triangle over the clip space like this:
// (the asterisks represent the clip space bounds)
//-1,1 1,1
// ---------------------------------
// | * .
// | * .
// | * .
// | * .
// | * .
// | * .
// |***************
// | . 1,-1
// | .
// | .
// | .
// | .
// |.
@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
var result: VertexOutput;
let x = i32(vertex_index) / 2;
let y = i32(vertex_index) & 1;
let tc = vec2<f32>(
f32(x) * 2.0,
f32(y) * 2.0
);
result.position = vec4<f32>(
tc.x * 2.0 - 1.0,
1.0 - tc.y * 2.0,
0.0, 1.0
);
result.tex_coords = tc;
return result;
}
struct Uniforms {
view_inv: mat4x4<f32>,
proj_inv: mat4x4<f32>,
};
@group(0) @binding(0)
var<uniform> uniforms: Uniforms;
@group(0) @binding(1)
var acc_struct: acceleration_structure;
@fragment
fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
var color = vec4<f32>(vertex.tex_coords, 0.0, 1.0);
let d = vertex.tex_coords * 2.0 - 1.0;
let origin = (uniforms.view_inv * vec4<f32>(0.0,0.0,0.0,1.0)).xyz;
let temp = uniforms.proj_inv * vec4<f32>(d.x, d.y, 1.0, 1.0);
let direction = (uniforms.view_inv * vec4<f32>(normalize(temp.xyz), 0.0)).xyz;
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
color = vec4<f32>(intersection.barycentrics, 1.0 - intersection.barycentrics.x - intersection.barycentrics.y, 1.0);
}
return color;
}

20
examples/src/ray_scene/cube.mtl Normal file
View File

@ -0,0 +1,20 @@
# Blender MTL File: 'None'
# Material Count: 2
newmtl Material
Ns 250.000000
Ka 1.000000 1.000000 1.000000
Kd 0.000000 0.009087 0.800000
Ks 0.500000 0.500000 0.500000
Ke 0.000000 0.000000 0.000000
Ni 1.450000
d 1.000000
illum 2
newmtl None
Ns 500
Ka 0.8 0.8 0.8
Kd 0.8 0.8 0.8
Ks 0.8 0.8 0.8
d 1
illum 2

2587
examples/src/ray_scene/cube.obj Normal file
View File

File diff suppressed because it is too large Load Diff

569
examples/src/ray_scene/mod.rs Normal file
View File

@ -0,0 +1,569 @@
use bytemuck::{Pod, Zeroable};
use glam::{Mat4, Quat, Vec3};
use std::f32::consts::PI;
use std::ops::IndexMut;
use std::{borrow::Cow, future::Future, iter, mem, ops::Range, pin::Pin, task, time::Instant};
use wgpu::util::DeviceExt;
// from cube
#[repr(C)]
#[derive(Debug, Clone, Copy, Pod, Zeroable, Default)]
struct Vertex {
pos: [f32; 3],
_p0: [u32; 1],
normal: [f32; 3],
_p1: [u32; 1],
uv: [f32; 2],
_p2: [u32; 2],
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Uniforms {
view_inverse: Mat4,
proj_inverse: Mat4,
}
/// A wrapper for `pop_error_scope` futures that panics if an error occurs.
///
/// Given a future `inner` of an `Option<E>` for some error type `E`,
/// wait for the future to be ready, and panic if its value is `Some`.
///
/// This can be done simpler with `FutureExt`, but we don't want to add
/// a dependency just for this small case.
struct ErrorFuture<F> {
inner: F,
}
impl<F: Future<Output = Option<wgpu::Error>>> Future for ErrorFuture<F> {
type Output = ();
fn poll(self: Pin<&mut Self>, cx: &mut task::Context<'_>) -> task::Poll<()> {
let inner = unsafe { self.map_unchecked_mut(|me| &mut me.inner) };
inner.poll(cx).map(|error| {
if let Some(e) = error {
panic!("Rendering {}", e);
}
})
}
}
#[derive(Debug, Clone, Default)]
struct RawSceneComponents {
vertices: Vec<Vertex>,
indices: Vec<u32>,
geometries: Vec<(Range<usize>, Material)>, // index range, material
instances: Vec<(Range<usize>, Range<usize>)>, // vertex range, geometry range
}
struct SceneComponents {
vertices: wgpu::Buffer,
indices: wgpu::Buffer,
geometries: wgpu::Buffer,
instances: wgpu::Buffer,
bottom_level_acceleration_structures: Vec<wgpu::Blas>,
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct InstanceEntry {
first_vertex: u32,
first_geometry: u32,
last_geometry: u32,
_pad: u32,
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable, Default)]
struct GeometryEntry {
first_index: u32,
_p0: [u32; 3],
material: Material,
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable, Default, Debug)]
struct Material {
roughness_exponent: f32,
metalness: f32,
specularity: f32,
_p0: [u32; 1],
albedo: [f32; 3],
_p1: [u32; 1],
}
fn load_model(scene: &mut RawSceneComponents, path: &str) {
let path = env!("CARGO_MANIFEST_DIR").to_string() + "/src" + path;
println!("{}", path);
let mut object = obj::Obj::load(path).unwrap();
object.load_mtls().unwrap();
let data = object.data;
let start_vertex_index = scene.vertices.len();
let start_geometry_index = scene.geometries.len();
let mut mapping = std::collections::HashMap::<(usize, usize, usize), usize>::new();
let mut next_index = 0;
for object in data.objects {
for group in object.groups {
let start_index_index = scene.indices.len();
for poly in group.polys {
for end_index in 2..poly.0.len() {
for &index in &[0, end_index - 1, end_index] {
let obj::IndexTuple(position_id, texture_id, normal_id) = poly.0[index];
let texture_id = texture_id.expect("uvs required");
let normal_id = normal_id.expect("normals required");
let index = *mapping
.entry((position_id, texture_id, normal_id))
.or_insert(next_index);
if index == next_index {
next_index += 1;
scene.vertices.push(Vertex {
pos: data.position[position_id],
uv: data.texture[texture_id],
normal: data.normal[normal_id],
..Default::default()
})
}
scene.indices.push(index as u32);
}
}
}
let mut material: Material = Default::default();
if let Some(obj::ObjMaterial::Mtl(mat)) = group.material {
if let Some(kd) = mat.kd {
material.albedo = kd;
}
if let Some(ns) = mat.ns {
material.roughness_exponent = ns;
}
if let Some(ka) = mat.ka {
material.metalness = ka[0];
}
if let Some(ks) = mat.ks {
material.specularity = ks[0];
}
}
scene
.geometries
.push((start_index_index..scene.indices.len(), material));
}
}
scene.instances.push((
start_vertex_index..scene.vertices.len(),
start_geometry_index..scene.geometries.len(),
));
// dbg!(scene.vertices.len());
// dbg!(scene.indices.len());
// dbg!(&scene.geometries);
// dbg!(&scene.instances);
}
fn upload_scene_components(
device: &wgpu::Device,
queue: &wgpu::Queue,
scene: &RawSceneComponents,
) -> SceneComponents {
let geometry_buffer_content = scene
.geometries
.iter()
.map(|(index_range, material)| GeometryEntry {
first_index: index_range.start as u32,
material: *material,
..Default::default()
})
.collect::<Vec<_>>();
let instance_buffer_content = scene
.instances
.iter()
.map(|geometry| InstanceEntry {
first_vertex: geometry.0.start as u32,
first_geometry: geometry.1.start as u32,
last_geometry: geometry.1.end as u32,
_pad: 1,
})
.collect::<Vec<_>>();
let vertices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Vertices"),
contents: bytemuck::cast_slice(&scene.vertices),
usage: wgpu::BufferUsages::VERTEX
| wgpu::BufferUsages::STORAGE
| wgpu::BufferUsages::BLAS_INPUT,
});
let indices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Indices"),
contents: bytemuck::cast_slice(&scene.indices),
usage: wgpu::BufferUsages::INDEX
| wgpu::BufferUsages::STORAGE
| wgpu::BufferUsages::BLAS_INPUT,
});
let geometries = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Geometries"),
contents: bytemuck::cast_slice(&geometry_buffer_content),
usage: wgpu::BufferUsages::STORAGE,
});
let instances = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Instances"),
contents: bytemuck::cast_slice(&instance_buffer_content),
usage: wgpu::BufferUsages::STORAGE,
});
let (size_descriptors, bottom_level_acceleration_structures): (Vec<_>, Vec<_>) = scene
.instances
.iter()
.map(|(vertex_range, geometry_range)| {
let size_desc: Vec<wgpu::BlasTriangleGeometrySizeDescriptor> = (*geometry_range)
.clone()
.map(|i| wgpu::BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
vertex_count: vertex_range.end as u32 - vertex_range.start as u32,
index_format: Some(wgpu::IndexFormat::Uint32),
index_count: Some(
scene.geometries[i].0.end as u32 - scene.geometries[i].0.start as u32,
),
flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
})
.collect();
let blas = device.create_blas(
&wgpu::CreateBlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
},
wgpu::BlasGeometrySizeDescriptors::Triangles {
descriptors: size_desc.clone(),
},
);
(size_desc, blas)
})
.unzip();
let build_entries: Vec<_> = scene
.instances
.iter()
.zip(size_descriptors.iter())
.zip(bottom_level_acceleration_structures.iter())
.map(|(((vertex_range, geometry_range), size_desc), blas)| {
let triangle_geometries: Vec<_> = size_desc
.iter()
.zip(geometry_range.clone())
.map(|(size, i)| wgpu::BlasTriangleGeometry {
size,
vertex_buffer: &vertices,
first_vertex: vertex_range.start as u32,
vertex_stride: mem::size_of::<Vertex>() as u64,
index_buffer: Some(&indices),
index_buffer_offset: Some(scene.geometries[i].0.start as u64 * 4),
transform_buffer: None,
transform_buffer_offset: None,
})
.collect();
wgpu::BlasBuildEntry {
blas,
geometry: wgpu::BlasGeometries::TriangleGeometries(triangle_geometries),
}
})
.collect();
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(build_entries.iter(), iter::empty());
queue.submit(Some(encoder.finish()));
SceneComponents {
vertices,
indices,
geometries,
instances,
bottom_level_acceleration_structures,
}
}
fn load_scene(device: &wgpu::Device, queue: &wgpu::Queue) -> SceneComponents {
let mut scene = RawSceneComponents::default();
load_model(&mut scene, "/skybox/models/teslacyberv3.0.obj");
load_model(&mut scene, "/ray_scene/cube.obj");
upload_scene_components(device, queue, &scene)
}
struct Example {
uniforms: Uniforms,
uniform_buf: wgpu::Buffer,
tlas_package: wgpu::TlasPackage,
pipeline: wgpu::RenderPipeline,
bind_group: wgpu::BindGroup,
start_inst: Instant,
scene_components: SceneComponents,
}
impl crate::framework::Example for Example {
fn required_features() -> wgpu::Features {
wgpu::Features::EXPERIMENTAL_RAY_QUERY
| wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
}
fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities {
wgpu::DownlevelCapabilities::default()
}
fn required_limits() -> wgpu::Limits {
wgpu::Limits::default()
}
fn init(
config: &wgpu::SurfaceConfiguration,
_adapter: &wgpu::Adapter,
device: &wgpu::Device,
queue: &wgpu::Queue,
) -> Self {
let side_count = 8;
let scene_components = load_scene(device, queue);
let uniforms = {
let view = Mat4::look_at_rh(Vec3::new(0.0, 0.0, 2.5), Vec3::ZERO, Vec3::Y);
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.001,
1000.0,
);
Uniforms {
view_inverse: view.inverse(),
proj_inverse: proj.inverse(),
}
};
let uniform_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Uniform Buffer"),
contents: bytemuck::cast_slice(&[uniforms]),
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
});
let tlas = device.create_tlas(&wgpu::CreateTlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
max_instances: side_count * side_count,
});
let tlas_package = wgpu::TlasPackage::new(tlas);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: None,
layout: None,
vertex: wgpu::VertexState {
module: &shader,
entry_point: Some("vs_main"),
compilation_options: Default::default(),
buffers: &[],
},
fragment: Some(wgpu::FragmentState {
module: &shader,
entry_point: Some("fs_main"),
compilation_options: Default::default(),
targets: &[Some(config.format.into())],
}),
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleList,
..Default::default()
},
depth_stencil: None,
multisample: wgpu::MultisampleState::default(),
multiview: None,
cache: None,
});
let bind_group_layout = pipeline.get_bind_group_layout(0);
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &bind_group_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: uniform_buf.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 5,
resource: tlas_package.as_binding(),
},
wgpu::BindGroupEntry {
binding: 1,
resource: scene_components.vertices.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 2,
resource: scene_components.indices.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 3,
resource: scene_components.geometries.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 4,
resource: scene_components.instances.as_entire_binding(),
},
],
});
let start_inst = Instant::now();
Example {
uniforms,
uniform_buf,
tlas_package,
pipeline,
bind_group,
start_inst,
scene_components,
}
}
fn update(&mut self, _event: winit::event::WindowEvent) {}
fn resize(
&mut self,
config: &wgpu::SurfaceConfiguration,
_device: &wgpu::Device,
queue: &wgpu::Queue,
) {
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.001,
1000.0,
);
self.uniforms.proj_inverse = proj.inverse();
queue.write_buffer(&self.uniform_buf, 0, bytemuck::cast_slice(&[self.uniforms]));
}
fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
device.push_error_scope(wgpu::ErrorFilter::Validation);
// scene update
{
let dist = 3.5;
let side_count = 2;
let anim_time = self.start_inst.elapsed().as_secs_f64() as f32;
for x in 0..side_count {
for y in 0..side_count {
let instance = self.tlas_package.index_mut(x + y * side_count);
let blas_index = (x + y)
% self
.scene_components
.bottom_level_acceleration_structures
.len();
let x = x as f32 / (side_count - 1) as f32;
let y = y as f32 / (side_count - 1) as f32;
let x = x * 2.0 - 1.0;
let y = y * 2.0 - 1.0;
let transform = Mat4::from_rotation_translation(
Quat::from_euler(
glam::EulerRot::XYZ,
anim_time * 0.5 * 0.342,
anim_time * 0.5 * 0.254,
anim_time * 0.5 * 0.832 + PI,
),
Vec3 {
x: x * dist,
y: y * dist,
z: -14.0,
},
);
let transform = transform.transpose().to_cols_array()[..12]
.try_into()
.unwrap();
*instance = Some(wgpu::TlasInstance::new(
&self.scene_components.bottom_level_acceleration_structures[blas_index],
transform,
blas_index as u32,
0xff,
));
}
}
}
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(iter::empty(), iter::once(&self.tlas_package));
{
let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view,
resolve_target: None,
ops: wgpu::Operations {
load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
store: wgpu::StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
rpass.set_pipeline(&self.pipeline);
rpass.set_bind_group(0, Some(&self.bind_group), &[]);
rpass.draw(0..3, 0..1);
}
queue.submit(Some(encoder.finish()));
}
}
pub fn main() {
crate::framework::run::<Example>("ray_scene");
}
#[cfg(test)]
#[wgpu_test::gpu_test]
static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
name: "ray_scene",
image_path: "/examples/src/ray_scene/screenshot.png",
width: 1024,
height: 768,
optional_features: wgpu::Features::default(),
base_test_parameters: wgpu_test::TestParameters {
required_features: <Example as crate::framework::Example>::required_features(),
required_limits: <Example as crate::framework::Example>::required_limits(),
force_fxc: false,
skips: vec![],
failures: Vec::new(),
required_downlevel_caps:
<Example as crate::framework::Example>::required_downlevel_capabilities(),
},
comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
_phantom: std::marker::PhantomData::<Example>,
};

BIN
examples/src/ray_scene/screenshot.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 76 KiB

164
examples/src/ray_scene/shader.wgsl Normal file
View File

@ -0,0 +1,164 @@
struct VertexOutput {
@builtin(position) position: vec4<f32>,
@location(0) tex_coords: vec2<f32>,
};
@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
var result: VertexOutput;
let x = i32(vertex_index) / 2;
let y = i32(vertex_index) & 1;
let tc = vec2<f32>(
f32(x) * 2.0,
f32(y) * 2.0
);
result.position = vec4<f32>(
tc.x * 2.0 - 1.0,
1.0 - tc.y * 2.0,
0.0, 1.0
);
result.tex_coords = tc;
return result;
}
/*
The contents of the RayQuery struct are roughly as follows
let RAY_FLAG_NONE = 0x00u;
let RAY_FLAG_OPAQUE = 0x01u;
let RAY_FLAG_NO_OPAQUE = 0x02u;
let RAY_FLAG_TERMINATE_ON_FIRST_HIT = 0x04u;
let RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08u;
let RAY_FLAG_CULL_BACK_FACING = 0x10u;
let RAY_FLAG_CULL_FRONT_FACING = 0x20u;
let RAY_FLAG_CULL_OPAQUE = 0x40u;
let RAY_FLAG_CULL_NO_OPAQUE = 0x80u;
let RAY_FLAG_SKIP_TRIANGLES = 0x100u;
let RAY_FLAG_SKIP_AABBS = 0x200u;
let RAY_QUERY_INTERSECTION_NONE = 0u;
let RAY_QUERY_INTERSECTION_TRIANGLE = 1u;
let RAY_QUERY_INTERSECTION_GENERATED = 2u;
let RAY_QUERY_INTERSECTION_AABB = 4u;
struct RayDesc {
flags: u32,
cull_mask: u32,
t_min: f32,
t_max: f32,
origin: vec3<f32>,
dir: vec3<f32>,
}
struct RayIntersection {
kind: u32,
t: f32,
instance_custom_index: u32,
instance_id: u32,
sbt_record_offset: u32,
geometry_index: u32,
primitive_index: u32,
barycentrics: vec2<f32>,
front_face: bool,
object_to_world: mat4x3<f32>,
world_to_object: mat4x3<f32>,
}
*/
struct Uniforms {
view_inv: mat4x4<f32>,
proj_inv: mat4x4<f32>,
};
struct Vertex {
pos: vec3<f32>,
normal: vec3<f32>,
uv: vec2<f32>,
};
struct Instance {
first_vertex: u32,
first_geometry: u32,
last_geometry: u32,
_pad: u32
};
struct Material{
roughness_exponent: f32,
metalness: f32,
specularity: f32,
albedo: vec3<f32>
}
struct Geometry {
first_index: u32,
material: Material,
};
@group(0) @binding(0)
var<uniform> uniforms: Uniforms;
@group(0) @binding(1)
var<storage, read> vertices: array<Vertex>;
@group(0) @binding(2)
var<storage, read> indices: array<u32>;
@group(0) @binding(3)
var<storage, read> geometries: array<Geometry>;
@group(0) @binding(4)
var<storage, read> instances: array<Instance>;
@group(0) @binding(5)
var acc_struct: acceleration_structure;
@fragment
fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
var color = vec4<f32>(vertex.tex_coords, 0.0, 1.0);
let d = vertex.tex_coords * 2.0 - 1.0;
let origin = (uniforms.view_inv * vec4<f32>(0.0,0.0,0.0,1.0)).xyz;
let temp = uniforms.proj_inv * vec4<f32>(d.x, d.y, 1.0, 1.0);
let direction = (uniforms.view_inv * vec4<f32>(normalize(temp.xyz), 0.0)).xyz;
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
let instance = instances[intersection.instance_custom_index];
let geometry = geometries[intersection.geometry_index + instance.first_geometry];
let index_offset = geometry.first_index;
let vertex_offset = instance.first_vertex;
let first_index_index = intersection.primitive_index * 3u + index_offset;
let v_0 = vertices[vertex_offset+indices[first_index_index+0u]];
let v_1 = vertices[vertex_offset+indices[first_index_index+1u]];
let v_2 = vertices[vertex_offset+indices[first_index_index+2u]];
let bary = vec3<f32>(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics);
let pos = v_0.pos * bary.x + v_1.pos * bary.y + v_2.pos * bary.z;
let normal_raw = v_0.normal * bary.x + v_1.normal * bary.y + v_2.normal * bary.z;
let uv = v_0.uv * bary.x + v_1.uv * bary.y + v_2.uv * bary.z;
let normal = normalize(normal_raw);
let material = geometry.material;
color = vec4<f32>(material.albedo, 1.0);
if(intersection.instance_custom_index == 1u){
color = vec4<f32>(normal, 1.0);
}
}
return color;
}

13
examples/src/ray_shadows/README.md Normal file
View File

@ -0,0 +1,13 @@
# ray-shadows
This example renders a ray traced shadow with hardware acceleration.
## To Run
```
cargo run --bin wgpu-examples ray_shadows
```
## Screenshots
![Shadow example](screenshot.png)

385
examples/src/ray_shadows/mod.rs Normal file
View File

@ -0,0 +1,385 @@
use std::{borrow::Cow, future::Future, iter, mem, pin::Pin, task, time::Instant};
use bytemuck::{Pod, Zeroable};
use glam::{Mat4, Vec3};
use wgpu::util::DeviceExt;
use wgpu::{vertex_attr_array, IndexFormat, VertexBufferLayout};
// from cube
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Vertex {
_pos: [f32; 3],
_normal: [f32; 3],
}
fn vertex(pos: [f32; 3], normal: [f32; 3]) -> Vertex {
Vertex {
_pos: pos,
_normal: normal,
}
}
fn create_vertices() -> (Vec<Vertex>, Vec<u16>) {
let vertex_data = [
// base
vertex([-1.0, 0.0, -1.0], [0.0, 1.0, 0.0]),
vertex([-1.0, 0.0, 1.0], [0.0, 1.0, 0.0]),
vertex([1.0, 0.0, -1.0], [0.0, 1.0, 0.0]),
vertex([1.0, 0.0, 1.0], [0.0, 1.0, 0.0]),
//shadow caster
vertex([-(1.0 / 3.0), 0.0, 1.0], [0.0, 0.0, 1.0]),
vertex([-(1.0 / 3.0), 2.0 / 3.0, 1.0], [0.0, 0.0, 1.0]),
vertex([1.0 / 3.0, 0.0, 1.0], [0.0, 0.0, 1.0]),
vertex([1.0 / 3.0, 2.0 / 3.0, 1.0], [0.0, 0.0, 1.0]),
];
let index_data: &[u16] = &[
0, 1, 2, 2, 3, 1, //base
4, 5, 6, 6, 7, 5,
];
(vertex_data.to_vec(), index_data.to_vec())
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Uniforms {
view_inverse: Mat4,
proj_inverse: Mat4,
vertex: Mat4,
}
/// A wrapper for `pop_error_scope` futures that panics if an error occurs.
///
/// Given a future `inner` of an `Option<E>` for some error type `E`,
/// wait for the future to be ready, and panic if its value is `Some`.
///
/// This can be done simpler with `FutureExt`, but we don't want to add
/// a dependency just for this small case.
struct ErrorFuture<F> {
inner: F,
}
impl<F: Future<Output = Option<wgpu::Error>>> Future for ErrorFuture<F> {
type Output = ();
fn poll(self: Pin<&mut Self>, cx: &mut task::Context<'_>) -> task::Poll<()> {
let inner = unsafe { self.map_unchecked_mut(|me| &mut me.inner) };
inner.poll(cx).map(|error| {
if let Some(e) = error {
panic!("Rendering {}", e);
}
})
}
}
struct Example {
uniforms: Uniforms,
uniform_buf: wgpu::Buffer,
vertex_buf: wgpu::Buffer,
index_buf: wgpu::Buffer,
pipeline: wgpu::RenderPipeline,
bind_group: wgpu::BindGroup,
start_inst: Instant,
}
const CAM_LOOK_AT: Vec3 = Vec3::new(0.0, 1.0, -1.5);
fn create_matrix(config: &wgpu::SurfaceConfiguration) -> Uniforms {
let view = Mat4::look_at_rh(CAM_LOOK_AT, Vec3::ZERO, Vec3::Y);
let proj = Mat4::perspective_rh(
59.0_f32.to_radians(),
config.width as f32 / config.height as f32,
0.1,
1000.0,
);
Uniforms {
view_inverse: view.inverse(),
proj_inverse: proj.inverse(),
vertex: (proj * view),
}
}
impl crate::framework::Example for Example {
fn required_features() -> wgpu::Features {
wgpu::Features::EXPERIMENTAL_RAY_QUERY
| wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
| wgpu::Features::PUSH_CONSTANTS
}
fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities {
wgpu::DownlevelCapabilities::default()
}
fn required_limits() -> wgpu::Limits {
wgpu::Limits {
max_push_constant_size: 12,
..wgpu::Limits::default()
}
}
fn init(
config: &wgpu::SurfaceConfiguration,
_adapter: &wgpu::Adapter,
device: &wgpu::Device,
queue: &wgpu::Queue,
) -> Self {
let uniforms = create_matrix(config);
let uniform_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Uniform Buffer"),
contents: bytemuck::cast_slice(&[uniforms]),
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
});
let (vertex_data, index_data) = create_vertices();
let vertex_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Vertex Buffer"),
contents: bytemuck::cast_slice(&vertex_data),
usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::BLAS_INPUT,
});
let index_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Index Buffer"),
contents: bytemuck::cast_slice(&index_data),
usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::BLAS_INPUT,
});
let blas_geo_size_desc = wgpu::BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
vertex_count: vertex_data.len() as u32,
index_format: Some(wgpu::IndexFormat::Uint16),
index_count: Some(index_data.len() as u32),
flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
};
let blas = device.create_blas(
&wgpu::CreateBlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
},
wgpu::BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_geo_size_desc.clone()],
},
);
let tlas = device.create_tlas(&wgpu::CreateTlasDescriptor {
label: None,
flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: wgpu::AccelerationStructureUpdateMode::Build,
max_instances: 1,
});
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: None,
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
});
let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: None,
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::FRAGMENT,
ty: wgpu::BindingType::AccelerationStructure,
count: None,
},
],
});
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: None,
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[wgpu::PushConstantRange {
stages: wgpu::ShaderStages::FRAGMENT,
range: 0..12,
}],
});
let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: None,
layout: Some(&pipeline_layout),
vertex: wgpu::VertexState {
module: &shader,
entry_point: Some("vs_main"),
compilation_options: Default::default(),
buffers: &[VertexBufferLayout {
array_stride: mem::size_of::<Vertex>() as wgpu::BufferAddress,
step_mode: Default::default(),
attributes: &vertex_attr_array![0 => Float32x3, 1 => Float32x3],
}],
},
fragment: Some(wgpu::FragmentState {
module: &shader,
entry_point: Some("fs_main"),
compilation_options: Default::default(),
targets: &[Some(config.format.into())],
}),
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleList,
..Default::default()
},
depth_stencil: None,
multisample: wgpu::MultisampleState::default(),
multiview: None,
cache: None,
});
let mut tlas_package = wgpu::TlasPackage::new(tlas);
tlas_package[0] = Some(wgpu::TlasInstance::new(
&blas,
[1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0],
0,
0xFF,
));
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(
iter::once(&wgpu::BlasBuildEntry {
blas: &blas,
geometry: wgpu::BlasGeometries::TriangleGeometries(vec![
wgpu::BlasTriangleGeometry {
size: &blas_geo_size_desc,
vertex_buffer: &vertex_buf,
first_vertex: 0,
vertex_stride: mem::size_of::<Vertex>() as u64,
index_buffer: Some(&index_buf),
index_buffer_offset: Some(0),
transform_buffer: None,
transform_buffer_offset: None,
},
]),
}),
iter::once(&tlas_package),
);
queue.submit(Some(encoder.finish()));
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &bind_group_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: uniform_buf.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 1,
resource: tlas_package.as_binding(),
},
],
});
let start_inst = Instant::now();
Example {
uniforms,
uniform_buf,
vertex_buf,
index_buf,
pipeline,
bind_group,
start_inst,
}
}
fn update(&mut self, _event: winit::event::WindowEvent) {}
fn resize(
&mut self,
config: &wgpu::SurfaceConfiguration,
_device: &wgpu::Device,
queue: &wgpu::Queue,
) {
self.uniforms = create_matrix(config);
queue.write_buffer(&self.uniform_buf, 0, bytemuck::cast_slice(&[self.uniforms]));
queue.submit(None);
}
fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
//device.push_error_scope(wgpu::ErrorFilter::Validation);
const LIGHT_DISTANCE: f32 = 5.0;
const TIME_SCALE: f32 = -0.2;
const INITIAL_TIME: f32 = 1.0;
let cos = (self.start_inst.elapsed().as_secs_f32() * TIME_SCALE + INITIAL_TIME).cos()
* LIGHT_DISTANCE;
let sin = (self.start_inst.elapsed().as_secs_f32() * TIME_SCALE + INITIAL_TIME).sin()
* LIGHT_DISTANCE;
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
{
let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view,
resolve_target: None,
ops: wgpu::Operations {
load: wgpu::LoadOp::Clear(wgpu::Color {
r: 0.1,
g: 0.1,
b: 0.1,
a: 1.0,
}),
store: wgpu::StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
rpass.set_pipeline(&self.pipeline);
rpass.set_bind_group(0, Some(&self.bind_group), &[]);
rpass.set_push_constants(wgpu::ShaderStages::FRAGMENT, 0, &0.0_f32.to_ne_bytes());
rpass.set_push_constants(wgpu::ShaderStages::FRAGMENT, 4, &cos.to_ne_bytes());
rpass.set_push_constants(wgpu::ShaderStages::FRAGMENT, 8, &sin.to_ne_bytes());
rpass.set_vertex_buffer(0, self.vertex_buf.slice(..));
rpass.set_index_buffer(self.index_buf.slice(..), IndexFormat::Uint16);
rpass.draw_indexed(0..12, 0, 0..1);
}
queue.submit(Some(encoder.finish()));
device.poll(wgpu::Maintain::Wait);
}
}
pub fn main() {
crate::framework::run::<Example>("ray-shadows");
}
#[cfg(test)]
#[wgpu_test::gpu_test]
static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
name: "ray_cube_shadows",
image_path: "/examples/src/ray_shadows/screenshot.png",
width: 1024,
height: 768,
optional_features: wgpu::Features::default(),
base_test_parameters: wgpu_test::TestParameters {
required_features: <Example as crate::framework::Example>::required_features(),
required_limits: <Example as crate::framework::Example>::required_limits(),
skips: vec![],
failures: Vec::new(),
required_downlevel_caps:
<Example as crate::framework::Example>::required_downlevel_capabilities(),
force_fxc: false,
},
comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
_phantom: std::marker::PhantomData::<Example>,
};

BIN
examples/src/ray_shadows/screenshot.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 12 KiB

70
examples/src/ray_shadows/shader.wgsl Normal file
View File

@ -0,0 +1,70 @@
struct VertexOutput {
@builtin(position) position: vec4<f32>,
@location(0) tex_coords: vec2<f32>,
@location(1) normal: vec3<f32>,
@location(2) world_position: vec3<f32>,
};
@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32, @location(0) position: vec3<f32>, @location(1) normal: vec3<f32>,) -> VertexOutput {
var result: VertexOutput;
let x = i32(vertex_index) / 2;
let y = i32(vertex_index) & 1;
let tc = vec2<f32>(
f32(x) * 2.0,
f32(y) * 2.0
);
result.tex_coords = tc;
result.position = uniforms.vertex * vec4<f32>(position, 1.0);
result.normal = normal;
result.world_position = position;
return result;
}
struct Uniforms {
view_inv: mat4x4<f32>,
proj_inv: mat4x4<f32>,
vertex: mat4x4<f32>,
};
@group(0) @binding(0)
var<uniform> uniforms: Uniforms;
@group(0) @binding(1)
var acc_struct: acceleration_structure;
var<push_constant> light: vec3<f32>;
const SURFACE_BRIGHTNESS = 0.5;
@fragment
fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
let camera = (uniforms.view_inv * vec4<f32>(0.0,0.0,0.0,1.0)).xyz;
var color = vec4<f32>(vertex.tex_coords, 0.0, 1.0);
let d = vertex.tex_coords * 2.0 - 1.0;
let origin = vertex.world_position;
let direction = normalize(light - vertex.world_position);
var normal: vec3<f32>;
let dir_cam = normalize(camera - vertex.world_position);
if (dot(dir_cam, vertex.normal) < 0.0) {
normal = -vertex.normal;
} else {
normal = vertex.normal;
}
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.0001, 200.0, origin, direction));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
color = vec4<f32>(vec3<f32>(0.1) * SURFACE_BRIGHTNESS, 1.0);
} else {
color = vec4<f32>(vec3<f32>(max(dot(direction, normal), 0.1)) * SURFACE_BRIGHTNESS, 1.0);
}
return color;
}

14
examples/src/ray_traced_triangle/README.md Normal file
View File

@ -0,0 +1,14 @@
# ray-traced-triangle
This example renders three triangles with hardware acceleration.
This is the same scene set-up as hal ray-traced triangle
## To Run
```
cargo run --bin wgpu-examples ray_traced_triangle
```
## Screenshots
![Triangle example](screenshot.png)

54
examples/src/ray_traced_triangle/blit.wgsl Normal file
View File

@ -0,0 +1,54 @@
// same as ray_cube_compute/blit.wgsl
struct VertexOutput {
@builtin(position) position: vec4<f32>,
@location(0) tex_coords: vec2<f32>,
};
// meant to be called with 3 vertex indices: 0, 1, 2
// draws one large triangle over the clip space like this:
// (the asterisks represent the clip space bounds)
//-1,1 1,1
// ---------------------------------
// | * .
// | * .
// | * .
// | * .
// | * .
// | * .
// |***************
// | . 1,-1
// | .
// | .
// | .
// | .
// |.
@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
var result: VertexOutput;
let x = i32(vertex_index) / 2;
let y = i32(vertex_index) & 1;
let tc = vec2<f32>(
f32(x) * 2.0,
f32(y) * 2.0
);
result.position = vec4<f32>(
tc.x * 2.0 - 1.0,
1.0 - tc.y * 2.0,
0.0, 1.0
);
result.tex_coords = tc;
return result;
}
@group(0)
@binding(0)
var r_color: texture_2d<f32>;
@group(0)
@binding(1)
var r_sampler: sampler;
@fragment
fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
return textureSample(r_color, r_sampler, vertex.tex_coords);
}

443
examples/src/ray_traced_triangle/mod.rs Normal file
View File

@ -0,0 +1,443 @@
use glam::{Mat4, Vec3};
use std::mem;
use std::time::Instant;
use wgpu::util::{BufferInitDescriptor, DeviceExt};
use wgpu::{include_wgsl, BufferUsages, IndexFormat, SamplerDescriptor};
use wgpu::{
AccelerationStructureFlags, AccelerationStructureUpdateMode, BlasBuildEntry, BlasGeometries,
BlasGeometrySizeDescriptors, BlasTriangleGeometry, BlasTriangleGeometrySizeDescriptor,
CreateBlasDescriptor, CreateTlasDescriptor, TlasInstance, TlasPackage,
};
struct Example {
tlas_package: TlasPackage,
compute_pipeline: wgpu::ComputePipeline,
blit_pipeline: wgpu::RenderPipeline,
bind_group: wgpu::BindGroup,
blit_bind_group: wgpu::BindGroup,
storage_texture: wgpu::Texture,
start: Instant,
}
#[repr(C)]
#[derive(bytemuck::Pod, bytemuck::Zeroable, Clone, Copy, Debug)]
struct Uniforms {
view_inverse: Mat4,
proj_inverse: Mat4,
}
impl crate::framework::Example for Example {
fn required_features() -> wgpu::Features {
wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
| wgpu::Features::EXPERIMENTAL_RAY_QUERY
}
fn required_limits() -> wgpu::Limits {
wgpu::Limits::default()
}
fn init(
config: &wgpu::SurfaceConfiguration,
_adapter: &wgpu::Adapter,
device: &wgpu::Device,
queue: &wgpu::Queue,
) -> Self {
let shader = device.create_shader_module(include_wgsl!("shader.wgsl"));
let blit_shader = device.create_shader_module(include_wgsl!("blit.wgsl"));
let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("bgl for shader.wgsl"),
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::StorageTexture {
access: wgpu::StorageTextureAccess::WriteOnly,
format: wgpu::TextureFormat::Rgba8Unorm,
view_dimension: wgpu::TextureViewDimension::D2,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 2,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::AccelerationStructure,
count: None,
},
],
});
let blit_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("bgl for blit.wgsl"),
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
ty: wgpu::BindingType::Texture {
sample_type: wgpu::TextureSampleType::Float { filterable: false },
view_dimension: wgpu::TextureViewDimension::D2,
multisampled: false,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::NonFiltering),
count: None,
},
],
});
let vertices: [f32; 9] = [1.0, 1.0, 0.0, -1.0, 1.0, 0.0, 0.0, -1.0, 0.0];
let indices: [u32; 3] = [0, 1, 2];
let vertex_buffer = device.create_buffer_init(&BufferInitDescriptor {
label: Some("vertex buffer"),
contents: bytemuck::cast_slice(&vertices),
usage: BufferUsages::BLAS_INPUT,
});
let index_buffer = device.create_buffer_init(&BufferInitDescriptor {
label: Some("vertex buffer"),
contents: bytemuck::cast_slice(&indices),
usage: BufferUsages::BLAS_INPUT,
});
let blas_size_desc = BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
// 3 coordinates per vertex
vertex_count: (vertices.len() / 3) as u32,
index_format: Some(IndexFormat::Uint32),
index_count: Some(indices.len() as u32),
flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
};
let blas = device.create_blas(
&CreateBlasDescriptor {
label: None,
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size_desc.clone()],
},
);
let tlas = device.create_tlas(&CreateTlasDescriptor {
label: None,
max_instances: 3,
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
});
let mut tlas_package = TlasPackage::new(tlas);
tlas_package[0] = Some(TlasInstance::new(
&blas,
Mat4::from_translation(Vec3 {
x: 0.0,
y: 0.0,
z: 0.0,
})
.transpose()
.to_cols_array()[..12]
.try_into()
.unwrap(),
0,
0xff,
));
tlas_package[1] = Some(TlasInstance::new(
&blas,
Mat4::from_translation(Vec3 {
x: -1.0,
y: -1.0,
z: -2.0,
})
.transpose()
.to_cols_array()[..12]
.try_into()
.unwrap(),
0,
0xff,
));
tlas_package[2] = Some(TlasInstance::new(
&blas,
Mat4::from_translation(Vec3 {
x: 1.0,
y: -1.0,
z: -2.0,
})
.transpose()
.to_cols_array()[..12]
.try_into()
.unwrap(),
0,
0xff,
));
let uniforms = {
let view = Mat4::look_at_rh(Vec3::new(0.0, 0.0, 2.5), Vec3::ZERO, Vec3::Y);
let proj = Mat4::perspective_rh(59.0_f32.to_radians(), 1.0, 0.001, 1000.0);
Uniforms {
view_inverse: view.inverse(),
proj_inverse: proj.inverse(),
}
};
let uniform_buffer = device.create_buffer_init(&BufferInitDescriptor {
label: None,
contents: bytemuck::cast_slice(&[uniforms]),
usage: BufferUsages::UNIFORM,
});
let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor::default());
encoder.build_acceleration_structures(
Some(&BlasBuildEntry {
blas: &blas,
geometry: BlasGeometries::TriangleGeometries(vec![BlasTriangleGeometry {
size: &blas_size_desc,
vertex_buffer: &vertex_buffer,
first_vertex: 0,
vertex_stride: mem::size_of::<[f32; 3]>() as wgpu::BufferAddress,
// in this case since one triangle gets no compression from an index buffer `index_buffer` and `index_buffer_offset` could be `None`.
index_buffer: Some(&index_buffer),
index_buffer_offset: Some(0),
transform_buffer: None,
transform_buffer_offset: None,
}]),
}),
Some(&tlas_package),
);
queue.submit(Some(encoder.finish()));
let storage_tex = device.create_texture(&wgpu::TextureDescriptor {
label: None,
size: wgpu::Extent3d {
width: config.width,
height: config.height,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba8Unorm,
usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
view_formats: &[],
});
let sampler = device.create_sampler(&SamplerDescriptor {
label: None,
address_mode_u: Default::default(),
address_mode_v: Default::default(),
address_mode_w: Default::default(),
mag_filter: wgpu::FilterMode::Nearest,
min_filter: wgpu::FilterMode::Nearest,
mipmap_filter: wgpu::FilterMode::Nearest,
lod_min_clamp: 1.0,
lod_max_clamp: 1.0,
compare: None,
anisotropy_clamp: 1,
border_color: None,
});
let compute_pipeline_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("pipeline layout for shader.wgsl"),
bind_group_layouts: &[&bgl],
push_constant_ranges: &[],
});
let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("pipeline for shader.wgsl"),
layout: Some(&compute_pipeline_layout),
module: &shader,
entry_point: None,
compilation_options: Default::default(),
cache: None,
});
let blit_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("pipeline layout for blit.wgsl"),
bind_group_layouts: &[&blit_bgl],
push_constant_ranges: &[],
});
let blit_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("pipeline for blit.wgsl"),
layout: Some(&blit_pipeline_layout),
vertex: wgpu::VertexState {
module: &blit_shader,
entry_point: None,
compilation_options: Default::default(),
buffers: &[],
},
primitive: Default::default(),
depth_stencil: None,
multisample: Default::default(),
fragment: Some(wgpu::FragmentState {
module: &blit_shader,
entry_point: None,
compilation_options: Default::default(),
targets: &[Some(wgpu::ColorTargetState {
format: config.format,
blend: None,
write_mask: Default::default(),
})],
}),
multiview: None,
cache: None,
});
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("bind group for shader.wgsl"),
layout: &bgl,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: uniform_buffer.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::TextureView(
&storage_tex.create_view(&wgpu::TextureViewDescriptor::default()),
),
},
wgpu::BindGroupEntry {
binding: 2,
resource: wgpu::BindingResource::AccelerationStructure(tlas_package.tlas()),
},
],
});
let blit_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("bind group for blit.wgsl"),
layout: &blit_bgl,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(
&storage_tex.create_view(&wgpu::TextureViewDescriptor::default()),
),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::Sampler(&sampler),
},
],
});
Self {
tlas_package,
compute_pipeline,
blit_pipeline,
bind_group,
blit_bind_group,
storage_texture: storage_tex,
start: Instant::now(),
}
}
fn resize(
&mut self,
_config: &wgpu::SurfaceConfiguration,
_device: &wgpu::Device,
_queue: &wgpu::Queue,
) {
}
fn update(&mut self, _event: winit::event::WindowEvent) {}
fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
self.tlas_package[0].as_mut().unwrap().transform =
Mat4::from_rotation_y(self.start.elapsed().as_secs_f32())
.transpose()
.to_cols_array()[..12]
.try_into()
.unwrap();
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
encoder.build_acceleration_structures(None, Some(&self.tlas_package));
{
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
cpass.set_pipeline(&self.compute_pipeline);
cpass.set_bind_group(0, Some(&self.bind_group), &[]);
cpass.dispatch_workgroups(
self.storage_texture.width() / 8,
self.storage_texture.height() / 8,
1,
);
}
{
let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view,
resolve_target: None,
ops: wgpu::Operations {
load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
store: wgpu::StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
rpass.set_pipeline(&self.blit_pipeline);
rpass.set_bind_group(0, Some(&self.blit_bind_group), &[]);
rpass.draw(0..3, 0..1);
}
queue.submit(Some(encoder.finish()));
}
}
pub fn main() {
crate::framework::run::<Example>("ray-traced-triangle");
}
#[cfg(test)]
#[wgpu_test::gpu_test]
static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
name: "ray_traced_triangle",
image_path: "/examples/src/ray_traced_triangle/screenshot.png",
width: 1024,
height: 768,
optional_features: wgpu::Features::default(),
base_test_parameters: wgpu_test::TestParameters {
required_features: <Example as crate::framework::Example>::required_features(),
required_limits: <Example as crate::framework::Example>::required_limits(),
force_fxc: false,
skips: vec![],
failures: Vec::new(),
required_downlevel_caps:
<Example as crate::framework::Example>::required_downlevel_capabilities(),
},
comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
_phantom: std::marker::PhantomData::<Example>,
};

BIN
examples/src/ray_traced_triangle/screenshot.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 116 KiB

39
examples/src/ray_traced_triangle/shader.wgsl Normal file
View File

@ -0,0 +1,39 @@
// duplicate of hal's ray-traced triangle shader
struct Uniforms {
view_inv: mat4x4<f32>,
proj_inv: mat4x4<f32>,
};
@group(0) @binding(0)
var<uniform> uniforms: Uniforms;
@group(0) @binding(1)
var output: texture_storage_2d<rgba8unorm, write>;
@group(0) @binding(2)
var acc_struct: acceleration_structure;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
let target_size = textureDimensions(output);
let pixel_center = vec2<f32>(global_id.xy) + vec2<f32>(0.5);
let in_uv = pixel_center / vec2<f32>(target_size.xy);
let d = in_uv * 2.0 - 1.0;
let origin = (uniforms.view_inv * vec4<f32>(0.0, 0.0, 0.0, 1.0)).xyz;
let temp = uniforms.proj_inv * vec4<f32>(d.x, d.y, 1.0, 1.0);
let direction = (uniforms.view_inv * vec4<f32>(normalize(temp.xyz), 0.0)).xyz;
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction));
rayQueryProceed(&rq);
var color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
let intersection = rayQueryGetCommittedIntersection(&rq);
if intersection.kind != RAY_QUERY_INTERSECTION_NONE {
color = vec4<f32>(intersection.barycentrics, 1.0 - intersection.barycentrics.x - intersection.barycentrics.y, 1.0);
}
textureStore(output, global_id.xy, color);
}

2
naga/Cargo.toml
View File

@ -77,7 +77,7 @@ indexmap.workspace = true
log = "0.4"
spirv = { version = "0.3", optional = true }
thiserror.workspace = true
serde = { version = "1.0.213", features = ["derive"], optional = true }
serde = { version = "1.0.214", 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 }

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

@ -1332,7 +1332,8 @@ impl<'a, W: Write> Writer<'a, W> {
crate::MathFunction::Pack4xI8
| crate::MathFunction::Pack4xU8
| crate::MathFunction::Unpack4xI8
| crate::MathFunction::Unpack4xU8 => {
| crate::MathFunction::Unpack4xU8
| crate::MathFunction::QuantizeToF16 => {
self.need_bake_expressions.insert(arg);
}
crate::MathFunction::ExtractBits => {
@ -3095,7 +3096,7 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_expr(image, ctx)?;
// All textureSize calls requires an lod argument
// except for multisampled samplers
if class.is_multisampled() {
if !class.is_multisampled() {
write!(self.out, ", 0")?;
}
write!(self.out, ")")?;
@ -3495,6 +3496,48 @@ impl<'a, W: Write> Writer<'a, W> {
Mf::Inverse => "inverse",
Mf::Transpose => "transpose",
Mf::Determinant => "determinant",
Mf::QuantizeToF16 => match *ctx.resolve_type(arg, &self.module.types) {
crate::TypeInner::Scalar { .. } => {
write!(self.out, "unpackHalf2x16(packHalf2x16(vec2(")?;
self.write_expr(arg, ctx)?;
write!(self.out, "))).x")?;
return Ok(());
}
crate::TypeInner::Vector {
size: crate::VectorSize::Bi,
..
} => {
write!(self.out, "unpackHalf2x16(packHalf2x16(")?;
self.write_expr(arg, ctx)?;
write!(self.out, "))")?;
return Ok(());
}
crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
..
} => {
write!(self.out, "vec3(unpackHalf2x16(packHalf2x16(")?;
self.write_expr(arg, ctx)?;
write!(self.out, ".xy)), unpackHalf2x16(packHalf2x16(")?;
self.write_expr(arg, ctx)?;
write!(self.out, ".zz)).x)")?;
return Ok(());
}
crate::TypeInner::Vector {
size: crate::VectorSize::Quad,
..
} => {
write!(self.out, "vec4(unpackHalf2x16(packHalf2x16(")?;
self.write_expr(arg, ctx)?;
write!(self.out, ".xy)), unpackHalf2x16(packHalf2x16(")?;
self.write_expr(arg, ctx)?;
write!(self.out, ".zw)))")?;
return Ok(());
}
_ => unreachable!(
"Correct TypeInner for QuantizeToF16 should be already validated"
),
},
// bits
Mf::CountTrailingZeros => {
match *ctx.resolve_type(arg, &self.module.types) {

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

@ -3036,6 +3036,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
Unpack4x8unorm,
Unpack4xI8,
Unpack4xU8,
QuantizeToF16,
Regular(&'static str),
MissingIntOverload(&'static str),
MissingIntReturnType(&'static str),
@ -3102,6 +3103,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
//Mf::Inverse =>,
Mf::Transpose => Function::Regular("transpose"),
Mf::Determinant => Function::Regular("determinant"),
Mf::QuantizeToF16 => Function::QuantizeToF16,
// bits
Mf::CountTrailingZeros => Function::CountTrailingZeros,
Mf::CountLeadingZeros => Function::CountLeadingZeros,
@ -3303,6 +3305,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
self.write_expr(module, arg, func_ctx)?;
write!(self.out, " >> 24) << 24 >> 24")?;
}
Function::QuantizeToF16 => {
write!(self.out, "f16tof32(f32tof16(")?;
self.write_expr(module, arg, func_ctx)?;
write!(self.out, "))")?;
}
Function::Regular(fun_name) => {
write!(self.out, "{fun_name}(")?;
self.write_expr(module, arg, func_ctx)?;

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

@ -1936,6 +1936,7 @@ impl<W: Write> Writer<W> {
Mf::Inverse => return Err(Error::UnsupportedCall(format!("{fun:?}"))),
Mf::Transpose => "transpose",
Mf::Determinant => "determinant",
Mf::QuantizeToF16 => "",
// bits
Mf::CountTrailingZeros => "ctz",
Mf::CountLeadingZeros => "clz",
@ -2144,6 +2145,22 @@ impl<W: Write> Writer<W> {
self.put_expression(arg, context, true)?;
write!(self.out, " >> 24) << 24 >> 24")?;
}
Mf::QuantizeToF16 => {
match *context.resolve_type(arg) {
crate::TypeInner::Scalar { .. } => write!(self.out, "float(half(")?,
crate::TypeInner::Vector { size, .. } => write!(
self.out,
"{NAMESPACE}::float{size}({NAMESPACE}::half{size}(",
size = back::vector_size_str(size),
)?,
_ => unreachable!(
"Correct TypeInner for QuantizeToF16 should be already validated"
),
};
self.put_expression(arg, context, true)?;
write!(self.out, "))")?;
}
_ => {
write!(self.out, "{NAMESPACE}::{fun_name}")?;
self.put_call_parameters(

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

@ -1032,6 +1032,12 @@ impl<'w> BlockContext<'w> {
arg0_id,
)),
Mf::Determinant => MathOp::Ext(spirv::GLOp::Determinant),
Mf::QuantizeToF16 => MathOp::Custom(Instruction::unary(
spirv::Op::QuantizeToF16,
result_type_id,
id,
arg0_id,
)),
Mf::ReverseBits => MathOp::Custom(Instruction::unary(
spirv::Op::BitReverse,
result_type_id,

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

@ -1723,6 +1723,7 @@ impl<W: Write> Writer<W> {
Mf::InverseSqrt => Function::Regular("inverseSqrt"),
Mf::Transpose => Function::Regular("transpose"),
Mf::Determinant => Function::Regular("determinant"),
Mf::QuantizeToF16 => Function::Regular("quantizeToF16"),
// bits
Mf::CountTrailingZeros => Function::Regular("countTrailingZeros"),
Mf::CountLeadingZeros => Function::Regular("countLeadingZeros"),

161
naga/src/diagnostic_filter.rs
View File

@ -1,9 +1,24 @@
//! [`DiagnosticFilter`]s and supporting functionality.
#[cfg(feature = "wgsl-in")]
use crate::Span;
use crate::{Arena, Handle};
#[cfg(feature = "arbitrary")]
use arbitrary::Arbitrary;
#[cfg(feature = "wgsl-in")]
use indexmap::IndexMap;
#[cfg(feature = "deserialize")]
use serde::Deserialize;
#[cfg(feature = "serialize")]
use serde::Serialize;
/// A severity set on a [`DiagnosticFilter`].
///
/// <https://www.w3.org/TR/WGSL/#diagnostic-severity>
#[derive(Clone, Copy, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Severity {
Off,
Info,
@ -33,7 +48,6 @@ impl Severity {
/// Naga does not yet support diagnostic items at lesser severities than
/// [`Severity::Error`]. When this is implemented, this method should be deleted, and the
/// severity should be used directly for reporting diagnostics.
#[cfg(feature = "wgsl-in")]
pub(crate) fn report_diag<E>(
self,
err: E,
@ -57,6 +71,9 @@ impl Severity {
///
/// <https://www.w3.org/TR/WGSL/#filterable-triggering-rules>
#[derive(Clone, Copy, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum FilterableTriggeringRule {
DerivativeUniformity,
}
@ -79,10 +96,13 @@ impl FilterableTriggeringRule {
}
}
#[cfg(feature = "wgsl-in")]
pub(crate) const fn tracking_issue_num(self) -> u16 {
/// The default severity associated with this triggering rule.
///
/// See <https://www.w3.org/TR/WGSL/#filterable-triggering-rules> for a table of default
/// severities.
pub(crate) const fn default_severity(self) -> Severity {
match self {
FilterableTriggeringRule::DerivativeUniformity => 5320,
FilterableTriggeringRule::DerivativeUniformity => Severity::Error,
}
}
}
@ -91,7 +111,140 @@ impl FilterableTriggeringRule {
///
/// <https://www.w3.org/TR/WGSL/#diagnostic-filter>
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct DiagnosticFilter {
pub new_severity: Severity,
pub triggering_rule: FilterableTriggeringRule,
}
/// A map of diagnostic filters to their severity and first occurrence's span.
///
/// Intended for front ends' first step into storing parsed [`DiagnosticFilter`]s.
#[derive(Clone, Debug, Default)]
#[cfg(feature = "wgsl-in")]
pub(crate) struct DiagnosticFilterMap(IndexMap<FilterableTriggeringRule, (Severity, Span)>);
#[cfg(feature = "wgsl-in")]
impl DiagnosticFilterMap {
pub(crate) fn new() -> Self {
Self::default()
}
/// Add the given `diagnostic_filter` parsed at the given `span` to this map.
pub(crate) fn add(
&mut self,
diagnostic_filter: DiagnosticFilter,
span: Span,
) -> Result<(), ConflictingDiagnosticRuleError> {
use indexmap::map::Entry;
let &mut Self(ref mut diagnostic_filters) = self;
let DiagnosticFilter {
new_severity,
triggering_rule,
} = diagnostic_filter;
match diagnostic_filters.entry(triggering_rule) {
Entry::Vacant(entry) => {
entry.insert((new_severity, span));
}
Entry::Occupied(entry) => {
let &(first_severity, first_span) = entry.get();
if first_severity != new_severity {
return Err(ConflictingDiagnosticRuleError {
triggering_rule,
triggering_rule_spans: [first_span, span],
});
}
}
}
Ok(())
}
}
#[cfg(feature = "wgsl-in")]
impl IntoIterator for DiagnosticFilterMap {
type Item = (FilterableTriggeringRule, (Severity, Span));
type IntoIter = indexmap::map::IntoIter<FilterableTriggeringRule, (Severity, Span)>;
fn into_iter(self) -> Self::IntoIter {
let Self(this) = self;
this.into_iter()
}
}
/// An error returned by [`DiagnosticFilterMap::add`] when it encounters conflicting rules.
#[cfg(feature = "wgsl-in")]
#[derive(Clone, Debug)]
pub(crate) struct ConflictingDiagnosticRuleError {
pub triggering_rule: FilterableTriggeringRule,
pub triggering_rule_spans: [Span; 2],
}
/// Represents a single parent-linking node in a tree of [`DiagnosticFilter`]s backed by a
/// [`crate::Arena`].
///
/// A single element of a _tree_ of diagnostic filter rules stored in
/// [`crate::Module::diagnostic_filters`]. When nodes are built by a front-end, module-applicable
/// filter rules are chained together in runs based on parse site. For instance, given the
/// following:
///
/// - Module-applicable rules `a` and `b`.
/// - Rules `c` and `d`, applicable to an entry point called `c_and_d_func`.
/// - Rule `e`, applicable to an entry point called `e_func`.
///
/// The tree would be represented as follows:
///
/// ```text
/// a <- b
/// ^
/// |- c <- d
/// |
/// \- e
/// ```
///
/// ...where:
///
/// - `d` is the first leaf consulted by validation in `c_and_d_func`.
/// - `e` is the first leaf consulted by validation in `e_func`.
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct DiagnosticFilterNode {
pub inner: DiagnosticFilter,
pub parent: Option<Handle<DiagnosticFilterNode>>,
}
impl DiagnosticFilterNode {
/// Finds the most specific filter rule applicable to `triggering_rule` from the chain of
/// diagnostic filter rules in `arena`, starting with `node`, and returns its severity. If none
/// is found, return the value of [`FilterableTriggeringRule::default_severity`].
///
/// When `triggering_rule` is not applicable to this node, its parent is consulted recursively.
pub(crate) fn search(
node: Option<Handle<Self>>,
arena: &Arena<Self>,
triggering_rule: FilterableTriggeringRule,
) -> Severity {
let mut next = node;
while let Some(handle) = next {
let node = &arena[handle];
let &Self { ref inner, parent } = node;
let &DiagnosticFilter {
triggering_rule: rule,
new_severity,
} = inner;
if rule == triggering_rule {
return new_severity;
}
next = parent;
}
triggering_rule.default_severity()
}
}

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

@ -1,4 +1,4 @@
use crate::diagnostic_filter::FilterableTriggeringRule;
use crate::diagnostic_filter::ConflictingDiagnosticRuleError;
use crate::front::wgsl::parse::directive::enable_extension::{
EnableExtension, UnimplementedEnableExtension,
};
@ -295,10 +295,13 @@ pub(crate) enum Error<'a> {
DiagnosticInvalidSeverity {
severity_control_name_span: Span,
},
DiagnosticNotYetImplemented {
triggering_rule: FilterableTriggeringRule,
span: Span,
},
DiagnosticDuplicateTriggeringRule(ConflictingDiagnosticRuleError),
}
impl<'a> From<ConflictingDiagnosticRuleError> for Error<'a> {
fn from(value: ConflictingDiagnosticRuleError) -> Self {
Self::DiagnosticDuplicateTriggeringRule(value)
}
}
#[derive(Clone, Debug)]
@ -1017,24 +1020,29 @@ impl<'a> Error<'a> {
)
.into()],
},
Error::DiagnosticNotYetImplemented {
Error::DiagnosticDuplicateTriggeringRule(ConflictingDiagnosticRuleError {
triggering_rule,
span,
} => ParseError {
message: format!(
"the `{}` diagnostic filter is not yet supported",
triggering_rule.to_ident()
),
labels: vec![(span, "".into())],
notes: vec![format!(
concat!(
"Let Naga maintainers know that you ran into this at ",
"<https://github.com/gfx-rs/wgpu/issues/{}>, ",
"so they can prioritize it!"
triggering_rule_spans,
}) => {
let [first_span, second_span] = triggering_rule_spans;
ParseError {
message: format!(
"found conflicting `diagnostic(…)` rule(s) for `{}`",
triggering_rule.to_ident()
),
triggering_rule.tracking_issue_num()
)],
},
labels: vec![
(first_span, "first rule".into()),
(second_span, "second rule".into()),
],
notes: vec![concat!(
"multiple `diagnostic(…)` rules with the same rule name ",
"conflict unless the severity is the same; ",
"delete the rule you don't want, or ",
"ensure that all severities with the same rule name match"
)
.into()],
}
}
}
}
}

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

@ -1013,7 +1013,11 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
&mut self,
tu: &'temp ast::TranslationUnit<'source>,
) -> Result<crate::Module, Error<'source>> {
let mut module = crate::Module::default();
let mut module = crate::Module {
diagnostic_filters: tu.diagnostic_filters.clone(),
diagnostic_filter_leaf: tu.diagnostic_filter_leaf,
..Default::default()
};
let mut ctx = GlobalContext {
ast_expressions: &tu.expressions,
@ -1244,7 +1248,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
.arguments
.iter()
.enumerate()
.map(|(i, arg)| {
.map(|(i, arg)| -> Result<_, Error<'_>> {
let ty = self.resolve_ast_type(arg.ty, ctx)?;
let expr = expressions
.append(crate::Expression::FunctionArgument(i as u32), arg.name.span);
@ -1263,7 +1267,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let result = f
.result
.as_ref()
.map(|res| {
.map(|res| -> Result<_, Error<'_>> {
let ty = self.resolve_ast_type(res.ty, ctx)?;
Ok(crate::FunctionResult {
ty,

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

@ -1,3 +1,4 @@
use crate::diagnostic_filter::DiagnosticFilterNode;
use crate::front::wgsl::parse::directive::enable_extension::EnableExtensions;
use crate::front::wgsl::parse::number::Number;
use crate::front::wgsl::Scalar;
@ -26,6 +27,17 @@ pub struct TranslationUnit<'a> {
/// These are referred to by `Handle<ast::Type<'a>>` values.
/// User-defined types are referred to by name until lowering.
pub types: Arena<Type<'a>>,
/// Arena for all diagnostic filter rules parsed in this module, including those in functions.
///
/// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
/// validation.
pub diagnostic_filters: Arena<DiagnosticFilterNode>,
/// The leaf of all `diagnostic(…)` directives in this module.
///
/// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
/// validation.
pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
}
#[derive(Debug, Clone, Copy)]

1
naga/src/front/wgsl/parse/conv.rs
View File

@ -230,6 +230,7 @@ pub fn map_standard_fun(word: &str) -> Option<crate::MathFunction> {
"inverseSqrt" => Mf::InverseSqrt,
"transpose" => Mf::Transpose,
"determinant" => Mf::Determinant,
"quantizeToF16" => Mf::QuantizeToF16,
// bits
"countTrailingZeros" => Mf::CountTrailingZeros,
"countLeadingZeros" => Mf::CountLeadingZeros,

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

@ -1,4 +1,6 @@
use crate::diagnostic_filter::{self, DiagnosticFilter, FilterableTriggeringRule};
use crate::diagnostic_filter::{
self, DiagnosticFilter, DiagnosticFilterMap, DiagnosticFilterNode, FilterableTriggeringRule,
};
use crate::front::wgsl::error::{Error, ExpectedToken};
use crate::front::wgsl::parse::directive::enable_extension::{
EnableExtension, EnableExtensions, UnimplementedEnableExtension,
@ -1907,10 +1909,8 @@ impl Parser {
let _ = lexer.next();
let mut body = ast::Block::default();
let (condition, span) = lexer.capture_span(|lexer| {
let condition = self.general_expression(lexer, ctx)?;
Ok(condition)
})?;
let (condition, span) =
lexer.capture_span(|lexer| self.general_expression(lexer, ctx))?;
let mut reject = ast::Block::default();
reject.stmts.push(ast::Statement {
kind: ast::StatementKind::Break,
@ -1966,11 +1966,12 @@ impl Parser {
let mut body = ast::Block::default();
if !lexer.skip(Token::Separator(';')) {
let (condition, span) = lexer.capture_span(|lexer| {
let condition = self.general_expression(lexer, ctx)?;
lexer.expect(Token::Separator(';'))?;
Ok(condition)
})?;
let (condition, span) =
lexer.capture_span(|lexer| -> Result<_, Error<'_>> {
let condition = self.general_expression(lexer, ctx)?;
lexer.expect(Token::Separator(';'))?;
Ok(condition)
})?;
let mut reject = ast::Block::default();
reject.stmts.push(ast::Statement {
kind: ast::StatementKind::Break,
@ -2523,6 +2524,7 @@ impl Parser {
let mut lexer = Lexer::new(source);
let mut tu = ast::TranslationUnit::default();
let mut enable_extensions = EnableExtensions::empty();
let mut diagnostic_filters = DiagnosticFilterMap::new();
// Parse directives.
while let Ok((ident, _directive_ident_span)) = lexer.peek_ident_with_span() {
@ -2532,12 +2534,8 @@ impl Parser {
match kind {
DirectiveKind::Diagnostic => {
if let Some(diagnostic_filter) = self.diagnostic_filter(&mut lexer)? {
let triggering_rule = diagnostic_filter.triggering_rule;
let span = self.peek_rule_span(&lexer);
Err(Error::DiagnosticNotYetImplemented {
triggering_rule,
span,
})?;
diagnostic_filters.add(diagnostic_filter, span)?;
}
lexer.expect(Token::Separator(';'))?;
}
@ -2583,6 +2581,8 @@ impl Parser {
lexer.enable_extensions = enable_extensions.clone();
tu.enable_extensions = enable_extensions;
tu.diagnostic_filter_leaf =
Self::write_diagnostic_filters(&mut tu.diagnostic_filters, diagnostic_filters, None);
loop {
match self.global_decl(&mut lexer, &mut tu) {
@ -2674,4 +2674,25 @@ impl Parser {
Ok(filter)
}
pub(crate) fn write_diagnostic_filters(
arena: &mut Arena<DiagnosticFilterNode>,
filters: DiagnosticFilterMap,
parent: Option<Handle<DiagnosticFilterNode>>,
) -> Option<Handle<DiagnosticFilterNode>> {
filters
.into_iter()
.fold(parent, |parent, (triggering_rule, (new_severity, span))| {
Some(arena.append(
DiagnosticFilterNode {
inner: DiagnosticFilter {
new_severity,
triggering_rule,
},
parent,
},
span,
))
})
}
}

15
naga/src/lib.rs
View File

@ -269,6 +269,7 @@ pub use crate::arena::{Arena, Handle, Range, UniqueArena};
pub use crate::span::{SourceLocation, Span, SpanContext, WithSpan};
#[cfg(feature = "arbitrary")]
use arbitrary::Arbitrary;
use diagnostic_filter::DiagnosticFilterNode;
#[cfg(feature = "deserialize")]
use serde::Deserialize;
#[cfg(feature = "serialize")]
@ -1199,6 +1200,7 @@ pub enum MathFunction {
Inverse,
Transpose,
Determinant,
QuantizeToF16,
// bits
CountTrailingZeros,
CountLeadingZeros,
@ -2286,4 +2288,17 @@ pub struct Module {
pub functions: Arena<Function>,
/// Entry points.
pub entry_points: Vec<EntryPoint>,
/// Arena for all diagnostic filter rules parsed in this module, including those in functions
/// and statements.
///
/// This arena contains elements of a _tree_ of diagnostic filter rules. When nodes are built
/// by a front-end, they refer to a parent scope
pub diagnostic_filters: Arena<DiagnosticFilterNode>,
/// The leaf of all diagnostic filter rules tree parsed from directives in this module.
///
/// In WGSL, this corresponds to `diagnostic(…);` directives.
///
/// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
/// validation.
pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
}

4
naga/src/proc/constant_evaluator.rs
View File

@ -137,8 +137,8 @@ macro_rules! gen_component_wise_extractor {
for idx in 0..(size as u8).into() {
let group = component_groups
.iter()
.map(|cs| cs[idx])
.collect::<ArrayVec<_, N>>()
.map(|cs| cs.get(idx).cloned().ok_or(err.clone()))
.collect::<Result<ArrayVec<_, N>, _>>()?
.into_inner()
.unwrap();
new_components.push($ident(

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

@ -478,6 +478,7 @@ impl super::MathFunction {
Self::Inverse => 1,
Self::Transpose => 1,
Self::Determinant => 1,
Self::QuantizeToF16 => 1,
// bits
Self::CountTrailingZeros => 1,
Self::CountLeadingZeros => 1,

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

@ -665,7 +665,8 @@ impl<'a> ResolveContext<'a> {
| Mf::Exp2
| Mf::Log
| Mf::Log2
| Mf::Pow => res_arg.clone(),
| Mf::Pow
| Mf::QuantizeToF16 => res_arg.clone(),
Mf::Modf | Mf::Frexp => {
let (size, width) = match res_arg.inner_with(types) {
&Ti::Scalar(crate::Scalar {

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

@ -6,6 +6,7 @@
//! - expression reference counts
use super::{ExpressionError, FunctionError, ModuleInfo, ShaderStages, ValidationFlags};
use crate::diagnostic_filter::{DiagnosticFilterNode, FilterableTriggeringRule};
use crate::span::{AddSpan as _, WithSpan};
use crate::{
arena::{Arena, Handle},
@ -15,10 +16,6 @@ use std::ops;
pub type NonUniformResult = Option<Handle<crate::Expression>>;
// Remove this once we update our uniformity analysis and
// add support for the `derivative_uniformity` diagnostic
const DISABLE_UNIFORMITY_REQ_FOR_FRAGMENT_STAGE: bool = true;
bitflags::bitflags! {
/// Kinds of expressions that require uniform control flow.
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
@ -26,8 +23,8 @@ bitflags::bitflags! {
#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub struct UniformityRequirements: u8 {
const WORK_GROUP_BARRIER = 0x1;
const DERIVATIVE = if DISABLE_UNIFORMITY_REQ_FOR_FRAGMENT_STAGE { 0 } else { 0x2 };
const IMPLICIT_LEVEL = if DISABLE_UNIFORMITY_REQ_FOR_FRAGMENT_STAGE { 0 } else { 0x4 };
const DERIVATIVE = 0x2;
const IMPLICIT_LEVEL = 0x4;
}
}
@ -289,6 +286,13 @@ pub struct FunctionInfo {
/// Indicates that the function is using dual source blending.
pub dual_source_blending: bool,
/// The leaf of all module-wide diagnostic filter rules tree parsed from directives in this
/// module.
///
/// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
/// validation.
diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
}
impl FunctionInfo {
@ -421,7 +425,10 @@ impl FunctionInfo {
let image_storage = match sampling.image {
GlobalOrArgument::Global(var) => GlobalOrArgument::Global(var),
GlobalOrArgument::Argument(i) => {
let handle = arguments[i as usize];
let Some(handle) = arguments.get(i as usize).cloned() else {
// Argument count mismatch, will be reported later by validate_call
break;
};
GlobalOrArgument::from_expression(expression_arena, handle).map_err(
|source| {
FunctionError::Expression { handle, source }
@ -434,7 +441,10 @@ impl FunctionInfo {
let sampler_storage = match sampling.sampler {
GlobalOrArgument::Global(var) => GlobalOrArgument::Global(var),
GlobalOrArgument::Argument(i) => {
let handle = arguments[i as usize];
let Some(handle) = arguments.get(i as usize).cloned() else {
// Argument count mismatch, will be reported later by validate_call
break;
};
GlobalOrArgument::from_expression(expression_arena, handle).map_err(
|source| {
FunctionError::Expression { handle, source }
@ -820,6 +830,7 @@ impl FunctionInfo {
other_functions: &[FunctionInfo],
mut disruptor: Option<UniformityDisruptor>,
expression_arena: &Arena<crate::Expression>,
diagnostic_filter_arena: &Arena<DiagnosticFilterNode>,
) -> Result<FunctionUniformity, WithSpan<FunctionError>> {
use crate::Statement as S;
@ -836,8 +847,21 @@ impl FunctionInfo {
&& !req.is_empty()
{
if let Some(cause) = disruptor {
return Err(FunctionError::NonUniformControlFlow(req, expr, cause)
.with_span_handle(expr, expression_arena));
let severity = DiagnosticFilterNode::search(
self.diagnostic_filter_leaf,
diagnostic_filter_arena,
FilterableTriggeringRule::DerivativeUniformity,
);
severity.report_diag(
FunctionError::NonUniformControlFlow(req, expr, cause)
.with_span_handle(expr, expression_arena),
// TODO: Yes, this isn't contextualized with source, because
// the user is supposed to render what would normally be an
// error here. Once we actually support warning-level
// diagnostic items, then we won't need this non-compliant hack:
// <https://github.com/gfx-rs/wgpu/issues/6458>
|e, level| log::log!(level, "{e}"),
)?;
}
}
requirements |= req;
@ -895,9 +919,13 @@ impl FunctionInfo {
exit: ExitFlags::empty(),
}
}
S::Block(ref b) => {
self.process_block(b, other_functions, disruptor, expression_arena)?
}
S::Block(ref b) => self.process_block(
b,
other_functions,
disruptor,
expression_arena,
diagnostic_filter_arena,
)?,
S::If {
condition,
ref accept,
@ -911,12 +939,14 @@ impl FunctionInfo {
other_functions,
branch_disruptor,
expression_arena,
diagnostic_filter_arena,
)?;
let reject_uniformity = self.process_block(
reject,
other_functions,
branch_disruptor,
expression_arena,
diagnostic_filter_arena,
)?;
accept_uniformity | reject_uniformity
}
@ -935,6 +965,7 @@ impl FunctionInfo {
other_functions,
case_disruptor,
expression_arena,
diagnostic_filter_arena,
)?;
case_disruptor = if case.fall_through {
case_disruptor.or(case_uniformity.exit_disruptor())
@ -950,14 +981,20 @@ impl FunctionInfo {
ref continuing,
break_if,
} => {
let body_uniformity =
self.process_block(body, other_functions, disruptor, expression_arena)?;
let body_uniformity = self.process_block(
body,
other_functions,
disruptor,
expression_arena,
diagnostic_filter_arena,
)?;
let continuing_disruptor = disruptor.or(body_uniformity.exit_disruptor());
let continuing_uniformity = self.process_block(
continuing,
other_functions,
continuing_disruptor,
expression_arena,
diagnostic_filter_arena,
)?;
if let Some(expr) = break_if {
let _ = self.add_ref(expr);
@ -1111,6 +1148,7 @@ impl ModuleInfo {
expressions: vec![ExpressionInfo::new(); fun.expressions.len()].into_boxed_slice(),
sampling: crate::FastHashSet::default(),
dual_source_blending: false,
diagnostic_filter_leaf: module.diagnostic_filter_leaf,
};
let resolve_context =
ResolveContext::with_locals(module, &fun.local_variables, &fun.arguments);
@ -1134,7 +1172,13 @@ impl ModuleInfo {
}
}
let uniformity = info.process_block(&fun.body, &self.functions, None, &fun.expressions)?;
let uniformity = info.process_block(
&fun.body,
&self.functions,
None,
&fun.expressions,
&module.diagnostic_filters,
)?;
info.uniformity = uniformity.result;
info.may_kill = uniformity.exit.contains(ExitFlags::MAY_KILL);
@ -1224,6 +1268,7 @@ fn uniform_control_flow() {
expressions: vec![ExpressionInfo::new(); expressions.len()].into_boxed_slice(),
sampling: crate::FastHashSet::default(),
dual_source_blending: false,
diagnostic_filter_leaf: None,
};
let resolve_context = ResolveContext {
constants: &Arena::new(),
@ -1270,7 +1315,8 @@ fn uniform_control_flow() {
&vec![stmt_emit1, stmt_if_uniform].into(),
&[],
None,
&expressions
&expressions,
&Arena::new(),
),
Ok(FunctionUniformity {
result: Uniformity {
@ -1297,26 +1343,59 @@ fn uniform_control_flow() {
reject: crate::Block::new(),
};
{
let block_info = info.process_block(
&vec![stmt_emit2.clone(), stmt_if_non_uniform.clone()].into(),
&[],
None,
&expressions,
&Arena::new(),
);
assert_eq!(
block_info,
Err(FunctionError::NonUniformControlFlow(
UniformityRequirements::DERIVATIVE,
derivative_expr,
UniformityDisruptor::Expression(non_uniform_global_expr)
)
.with_span()),
);
assert_eq!(info[derivative_expr].ref_count, 1);
// Test that the same thing passes when we disable the `derivative_uniformity`
let mut diagnostic_filters = Arena::new();
let diagnostic_filter_leaf = diagnostic_filters.append(
DiagnosticFilterNode {
inner: crate::diagnostic_filter::DiagnosticFilter {
new_severity: crate::diagnostic_filter::Severity::Off,
triggering_rule: FilterableTriggeringRule::DerivativeUniformity,
},
parent: None,
},
crate::Span::default(),
);
let mut info = FunctionInfo {
diagnostic_filter_leaf: Some(diagnostic_filter_leaf),
..info.clone()
};
let block_info = info.process_block(
&vec![stmt_emit2, stmt_if_non_uniform].into(),
&[],
None,
&expressions,
&diagnostic_filters,
);
if DISABLE_UNIFORMITY_REQ_FOR_FRAGMENT_STAGE {
assert_eq!(info[derivative_expr].ref_count, 2);
} else {
assert_eq!(
block_info,
Err(FunctionError::NonUniformControlFlow(
UniformityRequirements::DERIVATIVE,
derivative_expr,
UniformityDisruptor::Expression(non_uniform_global_expr)
)
.with_span()),
);
assert_eq!(info[derivative_expr].ref_count, 1);
}
assert_eq!(
block_info,
Ok(FunctionUniformity {
result: Uniformity {
non_uniform_result: None,
requirements: UniformityRequirements::DERIVATIVE,
},
exit: ExitFlags::empty()
}),
);
assert_eq!(info[derivative_expr].ref_count, 2);
}
assert_eq!(info[non_uniform_global], GlobalUse::READ);
@ -1329,7 +1408,8 @@ fn uniform_control_flow() {
&vec![stmt_emit3, stmt_return_non_uniform].into(),
&[],
Some(UniformityDisruptor::Return),
&expressions
&expressions,
&Arena::new(),
),
Ok(FunctionUniformity {
result: Uniformity {
@ -1356,7 +1436,8 @@ fn uniform_control_flow() {
&vec![stmt_emit4, stmt_assign, stmt_kill, stmt_return_pointer].into(),
&[],
Some(UniformityDisruptor::Discard),
&expressions
&expressions,
&Arena::new(),
),
Ok(FunctionUniformity {
result: Uniformity {

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

@ -39,12 +39,21 @@ pub enum ExpressionError {
IndexableLength(#[from] IndexableLengthError),
#[error("Operation {0:?} can't work with {1:?}")]
InvalidUnaryOperandType(crate::UnaryOperator, Handle<crate::Expression>),
#[error("Operation {0:?} can't work with {1:?} and {2:?}")]
InvalidBinaryOperandTypes(
crate::BinaryOperator,
Handle<crate::Expression>,
Handle<crate::Expression>,
),
#[error(
"Operation {:?} can't work with {:?} (of type {:?}) and {:?} (of type {:?})",
op,
lhs_expr,
lhs_type,
rhs_expr,
rhs_type
)]
InvalidBinaryOperandTypes {
op: crate::BinaryOperator,
lhs_expr: Handle<crate::Expression>,
lhs_type: crate::TypeInner,
rhs_expr: Handle<crate::Expression>,
rhs_type: crate::TypeInner,
},
#[error("Selecting is not possible")]
InvalidSelectTypes,
#[error("Relational argument {0:?} is not a boolean vector")]
@ -847,7 +856,13 @@ impl super::Validator {
function.expressions[right],
right_inner
);
return Err(ExpressionError::InvalidBinaryOperandTypes(op, left, right));
return Err(ExpressionError::InvalidBinaryOperandTypes {
op,
lhs_expr: left,
lhs_type: left_inner.clone(),
rhs_expr: right,
rhs_type: right_inner.clone(),
});
}
ShaderStages::all()
}
@ -1363,6 +1378,26 @@ impl super::Validator {
_ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)),
}
}
Mf::QuantizeToF16 => {
if arg1_ty.is_some() || arg2_ty.is_some() || arg3_ty.is_some() {
return Err(ExpressionError::WrongArgumentCount(fun));
}
match *arg_ty {
Ti::Scalar(Sc {
kind: Sk::Float,
width: 4,
})
| Ti::Vector {
scalar:
Sc {
kind: Sk::Float,
width: 4,
},
..
} => {}
_ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)),
}
}
// Remove once fixed https://github.com/gfx-rs/wgpu/issues/5276
Mf::CountLeadingZeros
| Mf::CountTrailingZeros

11
naga/src/valid/handles.rs
View File

@ -2,6 +2,7 @@
use crate::{
arena::{BadHandle, BadRangeError},
diagnostic_filter::DiagnosticFilterNode,
Handle,
};
@ -39,6 +40,8 @@ impl super::Validator {
ref types,
ref special_types,
ref global_expressions,
ref diagnostic_filters,
ref diagnostic_filter_leaf,
} = module;
// NOTE: Types being first is important. All other forms of validation depend on this.
@ -180,6 +183,14 @@ impl super::Validator {
validate_type(ty)?;
}
for (handle, _node) in diagnostic_filters.iter() {
let DiagnosticFilterNode { inner: _, parent } = diagnostic_filters[handle];
handle.check_dep_opt(parent)?;
}
if let Some(handle) = *diagnostic_filter_leaf {
handle.check_valid_for(diagnostic_filters)?;
}
Ok(())
}

2
naga/tests/in/diagnostic-filter.param.ron Normal file
View File

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

1
naga/tests/in/diagnostic-filter.wgsl Normal file
View File

@ -0,0 +1 @@
diagnostic(off, derivative_uniformity);

8
naga/tests/in/image.param.ron
View File

@ -3,5 +3,11 @@
version: (1, 1),
debug: true,
),
glsl_exclude_list: ["depth_load", "depth_no_comparison", "levels_queries"]
glsl: (
version: Desktop(430),
writer_flags: (""),
binding_map: {},
zero_initialize_workgroup_memory: true,
),
glsl_exclude_list: ["depth_load", "depth_no_comparison"]
)

3
naga/tests/in/image.wgsl
View File

@ -92,8 +92,9 @@ fn queries() -> @builtin(position) vec4<f32> {
@vertex
fn levels_queries() -> @builtin(position) vec4<f32> {
let num_levels_2d = textureNumLevels(image_2d);
let num_levels_2d_array = textureNumLevels(image_2d_array);
let num_layers_2d = textureNumLayers(image_2d_array);
let num_levels_2d_array = textureNumLevels(image_2d_array);
let num_layers_2d_array = textureNumLayers(image_2d_array);
let num_levels_cube = textureNumLevels(image_cube);
let num_levels_cube_array = textureNumLevels(image_cube_array);
let num_layers_cube = textureNumLayers(image_cube_array);

4
naga/tests/in/math-functions.wgsl
View File

@ -45,4 +45,8 @@ fn main() {
let frexp_b = frexp(1.5).fract;
let frexp_c: i32 = frexp(1.5).exp;
let frexp_d: i32 = frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp.x;
let quantizeToF16_a: f32 = quantizeToF16(1.0);
let quantizeToF16_b: vec2<f32> = quantizeToF16(vec2(1.0, 1.0));
let quantizeToF16_c: vec3<f32> = quantizeToF16(vec3(1.0, 1.0, 1.0));
let quantizeToF16_d: vec4<f32> = quantizeToF16(vec4(1.0, 1.0, 1.0, 1.0));
}

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

@ -1191,6 +1191,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2516,6 +2517,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2555,6 +2557,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2603,6 +2606,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2645,6 +2649,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2738,6 +2743,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2789,6 +2795,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2843,6 +2850,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2894,6 +2902,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -2948,6 +2957,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
entry_points: [
@ -3623,6 +3633,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -4074,6 +4085,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -4194,6 +4206,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -4257,6 +4270,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
const_expression_types: [

2
naga/tests/out/analysis/collatz.info.ron
View File

@ -274,6 +274,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
entry_points: [
@ -428,6 +429,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
const_expression_types: [],

1
naga/tests/out/analysis/overrides.info.ron
View File

@ -163,6 +163,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
const_expression_types: [

3
naga/tests/out/analysis/shadow.info.ron
View File

@ -412,6 +412,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
(
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@ -1571,6 +1572,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
entry_points: [
@ -1664,6 +1666,7 @@
],
sampling: [],
dual_source_blending: false,
diagnostic_filter_leaf: None,
),
],
const_expression_types: [

15
naga/tests/out/glsl/image.gather.Fragment.glsl
View File

@ -1,16 +1,11 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
#version 430 core
uniform sampler2D _group_0_binding_1_fs;
precision highp float;
precision highp int;
uniform usampler2D _group_0_binding_2_fs;
uniform highp sampler2D _group_0_binding_1_fs;
uniform isampler2D _group_0_binding_3_fs;
uniform highp usampler2D _group_0_binding_2_fs;
uniform highp isampler2D _group_0_binding_3_fs;
uniform highp sampler2DShadow _group_1_binding_2_fs;
uniform sampler2DShadow _group_1_binding_2_fs;
layout(location = 0) out vec4 _fs2p_location0;

30
naga/tests/out/glsl/image.levels_queries.Vertex.glsl Normal file
View File

@ -0,0 +1,30 @@
#version 430 core
#extension GL_ARB_shader_texture_image_samples : require
uniform sampler2D _group_0_binding_1_vs;
uniform sampler2DArray _group_0_binding_4_vs;
uniform samplerCube _group_0_binding_5_vs;
uniform samplerCubeArray _group_0_binding_6_vs;
uniform sampler3D _group_0_binding_7_vs;
uniform sampler2DMS _group_0_binding_8_vs;
void main() {
uint num_levels_2d = uint(textureQueryLevels(_group_0_binding_1_vs));
uint num_layers_2d = uint(textureSize(_group_0_binding_4_vs, 0).z);
uint num_levels_2d_array = uint(textureQueryLevels(_group_0_binding_4_vs));
uint num_layers_2d_array = uint(textureSize(_group_0_binding_4_vs, 0).z);
uint num_levels_cube = uint(textureQueryLevels(_group_0_binding_5_vs));
uint num_levels_cube_array = uint(textureQueryLevels(_group_0_binding_6_vs));
uint num_layers_cube = uint(textureSize(_group_0_binding_6_vs, 0).z);
uint num_levels_3d = uint(textureQueryLevels(_group_0_binding_7_vs));
uint num_samples_aa = uint(textureSamples(_group_0_binding_8_vs));
uint sum = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
gl_Position = vec4(float(sum));
return;
}

28
naga/tests/out/glsl/image.main.Compute.glsl
View File

@ -1,22 +1,18 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
precision highp float;
precision highp int;
#version 430 core
#extension GL_ARB_compute_shader : require
layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
uniform highp usampler2D _group_0_binding_0_cs;
uniform usampler2D _group_0_binding_0_cs;
uniform highp usampler2DMS _group_0_binding_3_cs;
uniform usampler2DMS _group_0_binding_3_cs;
layout(rgba8ui) readonly uniform highp uimage2D _group_0_binding_1_cs;
layout(rgba8ui) readonly uniform uimage2D _group_0_binding_1_cs;
uniform highp usampler2DArray _group_0_binding_5_cs;
uniform usampler2DArray _group_0_binding_5_cs;
uniform highp usampler2D _group_0_binding_7_cs;
uniform usampler1D _group_0_binding_7_cs;
layout(r32ui) writeonly uniform highp uimage2D _group_0_binding_2_cs;
layout(r32ui) writeonly uniform uimage1D _group_0_binding_2_cs;
void main() {
@ -28,15 +24,15 @@ void main() {
uvec4 value4_ = imageLoad(_group_0_binding_1_cs, itc);
uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, local_id.z), (int(local_id.z) + 1));
uvec4 value6_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
uvec4 value7_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0), int(local_id.z));
uvec4 value7_ = texelFetch(_group_0_binding_7_cs, int(local_id.x), int(local_id.z));
uvec4 value1u = texelFetch(_group_0_binding_0_cs, ivec2(uvec2(itc)), int(local_id.z));
uvec4 value2u = texelFetch(_group_0_binding_3_cs, ivec2(uvec2(itc)), int(local_id.z));
uvec4 value4u = imageLoad(_group_0_binding_1_cs, ivec2(uvec2(itc)));
uvec4 value5u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), local_id.z), (int(local_id.z) + 1));
uvec4 value6u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), int(local_id.z)), (int(local_id.z) + 1));
uvec4 value7u = texelFetch(_group_0_binding_7_cs, ivec2(uint(local_id.x), 0), int(local_id.z));
imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0), ((((value1_ + value2_) + value4_) + value5_) + value6_));
imageStore(_group_0_binding_2_cs, ivec2(uint(itc.x), 0), ((((value1u + value2u) + value4u) + value5u) + value6u));
uvec4 value7u = texelFetch(_group_0_binding_7_cs, int(uint(local_id.x)), int(local_id.z));
imageStore(_group_0_binding_2_cs, itc.x, ((((value1_ + value2_) + value4_) + value5_) + value6_));
imageStore(_group_0_binding_2_cs, int(uint(itc.x)), ((((value1u + value2u) + value4u) + value5u) + value6u));
return;
}

26
naga/tests/out/glsl/image.queries.Vertex.glsl
View File

@ -1,27 +1,22 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
#version 430 core
uniform sampler1D _group_0_binding_0_vs;
precision highp float;
precision highp int;
uniform sampler2D _group_0_binding_1_vs;
uniform highp sampler2D _group_0_binding_0_vs;
uniform sampler2DArray _group_0_binding_4_vs;
uniform highp sampler2D _group_0_binding_1_vs;
uniform samplerCube _group_0_binding_5_vs;
uniform highp sampler2DArray _group_0_binding_4_vs;
uniform samplerCubeArray _group_0_binding_6_vs;
uniform highp samplerCube _group_0_binding_5_vs;
uniform sampler3D _group_0_binding_7_vs;
uniform highp samplerCubeArray _group_0_binding_6_vs;
uniform highp sampler3D _group_0_binding_7_vs;
uniform highp sampler2DMS _group_0_binding_8_vs;
uniform sampler2DMS _group_0_binding_8_vs;
void main() {
uint dim_1d = uint(textureSize(_group_0_binding_0_vs, 0).x);
uint dim_1d_lod = uint(textureSize(_group_0_binding_0_vs, int(dim_1d)).x);
uint dim_1d = uint(textureSize(_group_0_binding_0_vs, 0));
uint dim_1d_lod = uint(textureSize(_group_0_binding_0_vs, int(dim_1d)));
uvec2 dim_2d = uvec2(textureSize(_group_0_binding_1_vs, 0).xy);
uvec2 dim_2d_lod = uvec2(textureSize(_group_0_binding_1_vs, 1).xy);
uvec2 dim_2d_array = uvec2(textureSize(_group_0_binding_4_vs, 0).xy);
@ -35,7 +30,6 @@ void main() {
uvec2 dim_2s_ms = uvec2(textureSize(_group_0_binding_8_vs).xy);
uint sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
gl_Position = vec4(float(sum));
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;
}

17
naga/tests/out/glsl/image.texture_sample.Fragment.glsl
View File

@ -1,16 +1,11 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
#version 430 core
uniform sampler1D _group_0_binding_0_fs;
precision highp float;
precision highp int;
uniform sampler2D _group_0_binding_1_fs;
uniform highp sampler2D _group_0_binding_0_fs;
uniform sampler2DArray _group_0_binding_4_fs;
uniform highp sampler2D _group_0_binding_1_fs;
uniform highp sampler2DArray _group_0_binding_4_fs;
uniform highp samplerCubeArray _group_0_binding_6_fs;
uniform samplerCubeArray _group_0_binding_6_fs;
layout(location = 0) out vec4 _fs2p_location0;
@ -18,7 +13,7 @@ void main() {
vec4 a = vec4(0.0);
vec2 tc = vec2(0.5);
vec3 tc3_ = vec3(0.5);
vec4 _e9 = texture(_group_0_binding_0_fs, vec2(tc.x, 0.0));
vec4 _e9 = texture(_group_0_binding_0_fs, tc.x);
vec4 _e10 = a;
a = (_e10 + _e9);
vec4 _e14 = texture(_group_0_binding_1_fs, vec2(tc));

13
naga/tests/out/glsl/image.texture_sample_comparison.Fragment.glsl
View File

@ -1,14 +1,9 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
#version 430 core
uniform sampler2DShadow _group_1_binding_2_fs;
precision highp float;
precision highp int;
uniform sampler2DArrayShadow _group_1_binding_3_fs;
uniform highp sampler2DShadow _group_1_binding_2_fs;
uniform highp sampler2DArrayShadow _group_1_binding_3_fs;
uniform highp samplerCubeShadow _group_1_binding_4_fs;
uniform samplerCubeShadow _group_1_binding_4_fs;
layout(location = 0) out float _fs2p_location0;

7
naga/tests/out/glsl/math-functions.main.Fragment.glsl
View File

@ -87,5 +87,12 @@ void main() {
float frexp_b = naga_frexp(1.5).fract_;
int frexp_c = naga_frexp(1.5).exp_;
int frexp_d = naga_frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp_.x;
float quantizeToF16_a = unpackHalf2x16(packHalf2x16(vec2(1.0))).x;
vec2 _e120 = vec2(1.0, 1.0);
vec2 quantizeToF16_b = unpackHalf2x16(packHalf2x16(_e120));
vec3 _e125 = vec3(1.0, 1.0, 1.0);
vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e125.xy)), unpackHalf2x16(packHalf2x16(_e125.zz)).x);
vec4 _e131 = vec4(1.0, 1.0, 1.0, 1.0);
vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e131.xy)), unpackHalf2x16(packHalf2x16(_e131.zw)));
}

7
naga/tests/out/hlsl/image.hlsl
View File

@ -177,14 +177,14 @@ uint NagaNumLevels2D(Texture2D<float4> tex)
return ret.z;
}
uint NagaNumLevels2DArray(Texture2DArray<float4> tex)
uint NagaNumLayers2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
uint NagaNumLayers2DArray(Texture2DArray<float4> tex)
uint NagaNumLevels2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
@ -229,8 +229,9 @@ uint NagaMSNumSamples2D(Texture2DMS<float4> tex)
float4 levels_queries() : SV_Position
{
uint num_levels_2d = NagaNumLevels2D(image_2d);
uint num_levels_2d_array = NagaNumLevels2DArray(image_2d_array);
uint num_layers_2d = NagaNumLayers2DArray(image_2d_array);
uint num_levels_2d_array = NagaNumLevels2DArray(image_2d_array);
uint num_layers_2d_array = NagaNumLayers2DArray(image_2d_array);
uint num_levels_cube = NagaNumLevelsCube(image_cube);
uint num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array);
uint num_layers_cube = NagaNumLayersCubeArray(image_cube_array);

4
naga/tests/out/hlsl/math-functions.hlsl
View File

@ -101,4 +101,8 @@ void main()
float frexp_b = naga_frexp(1.5).fract;
int frexp_c = naga_frexp(1.5).exp_;
int frexp_d = naga_frexp(float4(1.5, 1.5, 1.5, 1.5)).exp_.x;
float quantizeToF16_a = f16tof32(f32tof16(1.0));
float2 quantizeToF16_b = f16tof32(f32tof16(float2(1.0, 1.0)));
float3 quantizeToF16_c = f16tof32(f32tof16(float3(1.0, 1.0, 1.0)));
float4 quantizeToF16_d = f16tof32(f32tof16(float4(1.0, 1.0, 1.0, 1.0)));
}

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

@ -2476,4 +2476,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

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

@ -2476,4 +2476,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

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

@ -279,4 +279,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

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

@ -304,4 +304,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/collatz.compact.ron
View File

@ -330,4 +330,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/collatz.ron
View File

@ -330,4 +330,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/const_assert.compact.ron
View File

@ -51,4 +51,6 @@
),
],
entry_points: [],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/const_assert.ron
View File

@ -51,4 +51,6 @@
),
],
entry_points: [],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

24
naga/tests/out/ir/diagnostic-filter.compact.ron Normal file
View File

@ -0,0 +1,24 @@
(
types: [],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
overrides: [],
global_variables: [],
global_expressions: [],
functions: [],
entry_points: [],
diagnostic_filters: [
(
inner: (
new_severity: Off,
triggering_rule: DerivativeUniformity,
),
parent: None,
),
],
diagnostic_filter_leaf: Some(0),
)

24
naga/tests/out/ir/diagnostic-filter.ron Normal file
View File

@ -0,0 +1,24 @@
(
types: [],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
overrides: [],
global_variables: [],
global_expressions: [],
functions: [],
entry_points: [],
diagnostic_filters: [
(
inner: (
new_severity: Off,
triggering_rule: DerivativeUniformity,
),
parent: None,
),
],
diagnostic_filter_leaf: Some(0),
)

2
naga/tests/out/ir/fetch_depth.compact.ron
View File

@ -192,4 +192,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/fetch_depth.ron
View File

@ -262,4 +262,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/index-by-value.compact.ron
View File

@ -369,4 +369,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/index-by-value.ron
View File

@ -369,4 +369,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/local-const.compact.ron
View File

@ -136,4 +136,6 @@
),
],
entry_points: [],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/local-const.ron
View File

@ -136,4 +136,6 @@
),
],
entry_points: [],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
View File

@ -125,4 +125,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
View File

@ -125,4 +125,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides-ray-query.compact.ron
View File

@ -256,4 +256,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides-ray-query.ron
View File

@ -256,4 +256,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides.compact.ron
View File

@ -196,4 +196,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/overrides.ron
View File

@ -196,4 +196,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/shadow.compact.ron
View File

@ -1026,4 +1026,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/shadow.ron
View File

@ -1304,4 +1304,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/spec-constants.compact.ron
View File

@ -609,4 +609,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

2
naga/tests/out/ir/spec-constants.ron
View File

@ -715,4 +715,6 @@
),
),
],
diagnostic_filters: [],
diagnostic_filter_leaf: None,
)

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

@ -94,8 +94,9 @@ vertex levels_queriesOutput levels_queries(
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
uint num_levels_2d = image_2d.get_num_mip_levels();
uint num_levels_2d_array = image_2d_array.get_num_mip_levels();
uint num_layers_2d = image_2d_array.get_array_size();
uint num_levels_2d_array = image_2d_array.get_num_mip_levels();
uint num_layers_2d_array = image_2d_array.get_array_size();
uint num_levels_cube = image_cube.get_num_mip_levels();
uint num_levels_cube_array = image_cube_array.get_num_mip_levels();
uint num_layers_cube = image_cube_array.get_array_size();

4
naga/tests/out/msl/math-functions.msl
View File

@ -89,4 +89,8 @@ fragment void main_(
float frexp_b = naga_frexp(1.5).fract;
int frexp_c = naga_frexp(1.5).exp;
int frexp_d = naga_frexp(metal::float4(1.5, 1.5, 1.5, 1.5)).exp.x;
float quantizeToF16_a = float(half(1.0));
metal::float2 quantizeToF16_b = metal::float2(metal::half2(metal::float2(1.0, 1.0)));
metal::float3 quantizeToF16_c = metal::float3(metal::half3(metal::float3(1.0, 1.0, 1.0)));
metal::float4 quantizeToF16_d = metal::float4(metal::half4(metal::float4(1.0, 1.0, 1.0, 1.0)));
}

654
naga/tests/out/spv/image.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 518
; Bound: 520
OpCapability Shader
OpCapability Image1D
OpCapability Sampled1D
@ -13,16 +13,16 @@ OpEntryPoint GLCompute %78 "main" %75
OpEntryPoint GLCompute %169 "depth_load" %167
OpEntryPoint Vertex %189 "queries" %187
OpEntryPoint Vertex %241 "levels_queries" %240
OpEntryPoint Fragment %270 "texture_sample" %269
OpEntryPoint Fragment %417 "texture_sample_comparison" %415
OpEntryPoint Fragment %473 "gather" %472
OpEntryPoint Fragment %507 "depth_no_comparison" %506
OpEntryPoint Fragment %272 "texture_sample" %271
OpEntryPoint Fragment %419 "texture_sample_comparison" %417
OpEntryPoint Fragment %475 "gather" %474
OpEntryPoint Fragment %509 "depth_no_comparison" %508
OpExecutionMode %78 LocalSize 16 1 1
OpExecutionMode %169 LocalSize 16 1 1
OpExecutionMode %270 OriginUpperLeft
OpExecutionMode %417 OriginUpperLeft
OpExecutionMode %473 OriginUpperLeft
OpExecutionMode %507 OriginUpperLeft
OpExecutionMode %272 OriginUpperLeft
OpExecutionMode %419 OriginUpperLeft
OpExecutionMode %475 OriginUpperLeft
OpExecutionMode %509 OriginUpperLeft
OpName %31 "image_mipmapped_src"
OpName %33 "image_multisampled_src"
OpName %35 "image_depth_multisampled_src"
@ -51,12 +51,12 @@ OpName %167 "local_id"
OpName %169 "depth_load"
OpName %189 "queries"
OpName %241 "levels_queries"
OpName %270 "texture_sample"
OpName %284 "a"
OpName %417 "texture_sample_comparison"
OpName %422 "a"
OpName %473 "gather"
OpName %507 "depth_no_comparison"
OpName %272 "texture_sample"
OpName %286 "a"
OpName %419 "texture_sample_comparison"
OpName %424 "a"
OpName %475 "gather"
OpName %509 "depth_no_comparison"
OpDecorate %31 DescriptorSet 0
OpDecorate %31 Binding 0
OpDecorate %33 DescriptorSet 0
@ -108,10 +108,10 @@ OpDecorate %75 BuiltIn LocalInvocationId
OpDecorate %167 BuiltIn LocalInvocationId
OpDecorate %187 BuiltIn Position
OpDecorate %240 BuiltIn Position
OpDecorate %269 Location 0
OpDecorate %415 Location 0
OpDecorate %472 Location 0
OpDecorate %506 Location 0
OpDecorate %271 Location 0
OpDecorate %417 Location 0
OpDecorate %474 Location 0
OpDecorate %508 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 1 Unknown
@ -198,36 +198,36 @@ OpDecorate %506 Location 0
%187 = OpVariable %188 Output
%198 = OpConstant %4 0
%240 = OpVariable %188 Output
%269 = OpVariable %188 Output
%276 = OpConstant %7 0.5
%277 = OpTypeVector %7 2
%278 = OpConstantComposite %277 %276 %276
%279 = OpTypeVector %7 3
%280 = OpConstantComposite %279 %276 %276 %276
%281 = OpConstant %7 2.3
%282 = OpConstant %7 2.0
%283 = OpConstant %14 0
%285 = OpTypePointer Function %23
%286 = OpConstantNull %23
%289 = OpTypeSampledImage %15
%294 = OpTypeSampledImage %16
%315 = OpTypeSampledImage %18
%376 = OpTypeSampledImage %20
%416 = OpTypePointer Output %7
%415 = OpVariable %416 Output
%423 = OpTypePointer Function %7
%424 = OpConstantNull %7
%426 = OpTypeSampledImage %25
%431 = OpTypeSampledImage %26
%444 = OpTypeSampledImage %27
%451 = OpConstant %7 0.0
%472 = OpVariable %188 Output
%483 = OpConstant %4 1
%486 = OpConstant %4 3
%491 = OpTypeSampledImage %3
%494 = OpTypeVector %14 4
%495 = OpTypeSampledImage %17
%506 = OpVariable %188 Output
%271 = OpVariable %188 Output
%278 = OpConstant %7 0.5
%279 = OpTypeVector %7 2
%280 = OpConstantComposite %279 %278 %278
%281 = OpTypeVector %7 3
%282 = OpConstantComposite %281 %278 %278 %278
%283 = OpConstant %7 2.3
%284 = OpConstant %7 2.0
%285 = OpConstant %14 0
%287 = OpTypePointer Function %23
%288 = OpConstantNull %23
%291 = OpTypeSampledImage %15
%296 = OpTypeSampledImage %16
%317 = OpTypeSampledImage %18
%378 = OpTypeSampledImage %20
%418 = OpTypePointer Output %7
%417 = OpVariable %418 Output
%425 = OpTypePointer Function %7
%426 = OpConstantNull %7
%428 = OpTypeSampledImage %25
%433 = OpTypeSampledImage %26
%446 = OpTypeSampledImage %27
%453 = OpConstant %7 0.0
%474 = OpVariable %188 Output
%485 = OpConstant %4 1
%488 = OpConstant %4 3
%493 = OpTypeSampledImage %3
%496 = OpTypeVector %14 4
%497 = OpTypeSampledImage %17
%508 = OpVariable %188 Output
%78 = OpFunction %2 None %79
%74 = OpLabel
%77 = OpLoad %12 %75
@ -403,290 +403,292 @@ OpFunctionEnd
OpBranch %248
%248 = OpLabel
%249 = OpImageQueryLevels %4 %242
%250 = OpImageQueryLevels %4 %243
%251 = OpImageQuerySizeLod %12 %243 %198
%252 = OpCompositeExtract %4 %251 2
%253 = OpImageQueryLevels %4 %244
%254 = OpImageQueryLevels %4 %245
%255 = OpImageQuerySizeLod %12 %245 %198
%256 = OpCompositeExtract %4 %255 2
%257 = OpImageQueryLevels %4 %246
%258 = OpImageQuerySamples %4 %247
%259 = OpIAdd %4 %252 %256
%260 = OpIAdd %4 %259 %258
%261 = OpIAdd %4 %260 %249
%262 = OpIAdd %4 %261 %250
%263 = OpIAdd %4 %262 %257
%264 = OpIAdd %4 %263 %253
%265 = OpIAdd %4 %264 %254
%266 = OpConvertUToF %7 %265
%267 = OpCompositeConstruct %23 %266 %266 %266 %266
OpStore %240 %267
%250 = OpImageQuerySizeLod %12 %243 %198
%251 = OpCompositeExtract %4 %250 2
%252 = OpImageQueryLevels %4 %243
%253 = OpImageQuerySizeLod %12 %243 %198
%254 = OpCompositeExtract %4 %253 2
%255 = OpImageQueryLevels %4 %244
%256 = OpImageQueryLevels %4 %245
%257 = OpImageQuerySizeLod %12 %245 %198
%258 = OpCompositeExtract %4 %257 2
%259 = OpImageQueryLevels %4 %246
%260 = OpImageQuerySamples %4 %247
%261 = OpIAdd %4 %251 %258
%262 = OpIAdd %4 %261 %260
%263 = OpIAdd %4 %262 %249
%264 = OpIAdd %4 %263 %252
%265 = OpIAdd %4 %264 %259
%266 = OpIAdd %4 %265 %255
%267 = OpIAdd %4 %266 %256
%268 = OpConvertUToF %7 %267
%269 = OpCompositeConstruct %23 %268 %268 %268 %268
OpStore %240 %269
OpReturn
OpFunctionEnd
%270 = OpFunction %2 None %79
%268 = OpLabel
%284 = OpVariable %285 Function %286
%271 = OpLoad %15 %47
%272 = OpLoad %16 %49
%273 = OpLoad %18 %54
%274 = OpLoad %20 %58
%275 = OpLoad %24 %64
OpBranch %287
%287 = OpLabel
%288 = OpCompositeExtract %7 %278 0
%290 = OpSampledImage %289 %271 %275
%291 = OpImageSampleImplicitLod %23 %290 %288
%292 = OpLoad %23 %284
%293 = OpFAdd %23 %292 %291
OpStore %284 %293
%295 = OpSampledImage %294 %272 %275
%296 = OpImageSampleImplicitLod %23 %295 %278
%297 = OpLoad %23 %284
%298 = OpFAdd %23 %297 %296
OpStore %284 %298
%299 = OpSampledImage %294 %272 %275
%300 = OpImageSampleImplicitLod %23 %299 %278 ConstOffset %30
%301 = OpLoad %23 %284
%302 = OpFAdd %23 %301 %300
OpStore %284 %302
%303 = OpSampledImage %294 %272 %275
%304 = OpImageSampleExplicitLod %23 %303 %278 Lod %281
%305 = OpLoad %23 %284
%306 = OpFAdd %23 %305 %304
OpStore %284 %306
%307 = OpSampledImage %294 %272 %275
%308 = OpImageSampleExplicitLod %23 %307 %278 Lod|ConstOffset %281 %30
%309 = OpLoad %23 %284
%310 = OpFAdd %23 %309 %308
OpStore %284 %310
%311 = OpSampledImage %294 %272 %275
%312 = OpImageSampleImplicitLod %23 %311 %278 Bias|ConstOffset %282 %30
%313 = OpLoad %23 %284
%314 = OpFAdd %23 %313 %312
OpStore %284 %314
%316 = OpConvertUToF %7 %198
%317 = OpCompositeConstruct %279 %278 %316
%318 = OpSampledImage %315 %273 %275
%319 = OpImageSampleImplicitLod %23 %318 %317
%320 = OpLoad %23 %284
%321 = OpFAdd %23 %320 %319
OpStore %284 %321
%322 = OpConvertUToF %7 %198
%323 = OpCompositeConstruct %279 %278 %322
%324 = OpSampledImage %315 %273 %275
%325 = OpImageSampleImplicitLod %23 %324 %323 ConstOffset %30
%326 = OpLoad %23 %284
%327 = OpFAdd %23 %326 %325
OpStore %284 %327
%328 = OpConvertUToF %7 %198
%329 = OpCompositeConstruct %279 %278 %328
%330 = OpSampledImage %315 %273 %275
%331 = OpImageSampleExplicitLod %23 %330 %329 Lod %281
%332 = OpLoad %23 %284
%333 = OpFAdd %23 %332 %331
OpStore %284 %333
%334 = OpConvertUToF %7 %198
%335 = OpCompositeConstruct %279 %278 %334
%336 = OpSampledImage %315 %273 %275
%337 = OpImageSampleExplicitLod %23 %336 %335 Lod|ConstOffset %281 %30
%338 = OpLoad %23 %284
%339 = OpFAdd %23 %338 %337
OpStore %284 %339
%340 = OpConvertUToF %7 %198
%341 = OpCompositeConstruct %279 %278 %340
%342 = OpSampledImage %315 %273 %275
%343 = OpImageSampleImplicitLod %23 %342 %341 Bias|ConstOffset %282 %30
%344 = OpLoad %23 %284
%345 = OpFAdd %23 %344 %343
OpStore %284 %345
%346 = OpConvertSToF %7 %283
%347 = OpCompositeConstruct %279 %278 %346
%348 = OpSampledImage %315 %273 %275
%349 = OpImageSampleImplicitLod %23 %348 %347
%350 = OpLoad %23 %284
%351 = OpFAdd %23 %350 %349
OpStore %284 %351
%352 = OpConvertSToF %7 %283
%353 = OpCompositeConstruct %279 %278 %352
%354 = OpSampledImage %315 %273 %275
%355 = OpImageSampleImplicitLod %23 %354 %353 ConstOffset %30
%356 = OpLoad %23 %284
%357 = OpFAdd %23 %356 %355
OpStore %284 %357
%358 = OpConvertSToF %7 %283
%359 = OpCompositeConstruct %279 %278 %358
%360 = OpSampledImage %315 %273 %275
%361 = OpImageSampleExplicitLod %23 %360 %359 Lod %281
%362 = OpLoad %23 %284
%363 = OpFAdd %23 %362 %361
OpStore %284 %363
%364 = OpConvertSToF %7 %283
%365 = OpCompositeConstruct %279 %278 %364
%366 = OpSampledImage %315 %273 %275
%367 = OpImageSampleExplicitLod %23 %366 %365 Lod|ConstOffset %281 %30
%368 = OpLoad %23 %284
%369 = OpFAdd %23 %368 %367
OpStore %284 %369
%370 = OpConvertSToF %7 %283
%371 = OpCompositeConstruct %279 %278 %370
%372 = OpSampledImage %315 %273 %275
%373 = OpImageSampleImplicitLod %23 %372 %371 Bias|ConstOffset %282 %30
%374 = OpLoad %23 %284
%375 = OpFAdd %23 %374 %373
OpStore %284 %375
%377 = OpConvertUToF %7 %198
%378 = OpCompositeConstruct %23 %280 %377
%379 = OpSampledImage %376 %274 %275
%380 = OpImageSampleImplicitLod %23 %379 %378
%381 = OpLoad %23 %284
%382 = OpFAdd %23 %381 %380
OpStore %284 %382
%383 = OpConvertUToF %7 %198
%384 = OpCompositeConstruct %23 %280 %383
%385 = OpSampledImage %376 %274 %275
%386 = OpImageSampleExplicitLod %23 %385 %384 Lod %281
%387 = OpLoad %23 %284
%388 = OpFAdd %23 %387 %386
OpStore %284 %388
%389 = OpConvertUToF %7 %198
%390 = OpCompositeConstruct %23 %280 %389
%391 = OpSampledImage %376 %274 %275
%392 = OpImageSampleImplicitLod %23 %391 %390 Bias %282
%393 = OpLoad %23 %284
%394 = OpFAdd %23 %393 %392
OpStore %284 %394
%395 = OpConvertSToF %7 %283
%396 = OpCompositeConstruct %23 %280 %395
%397 = OpSampledImage %376 %274 %275
%398 = OpImageSampleImplicitLod %23 %397 %396
%399 = OpLoad %23 %284
%400 = OpFAdd %23 %399 %398
OpStore %284 %400
%401 = OpConvertSToF %7 %283
%402 = OpCompositeConstruct %23 %280 %401
%403 = OpSampledImage %376 %274 %275
%404 = OpImageSampleExplicitLod %23 %403 %402 Lod %281
%405 = OpLoad %23 %284
%406 = OpFAdd %23 %405 %404
OpStore %284 %406
%407 = OpConvertSToF %7 %283
%408 = OpCompositeConstruct %23 %280 %407
%409 = OpSampledImage %376 %274 %275
%410 = OpImageSampleImplicitLod %23 %409 %408 Bias %282
%411 = OpLoad %23 %284
%412 = OpFAdd %23 %411 %410
OpStore %284 %412
%413 = OpLoad %23 %284
OpStore %269 %413
%272 = OpFunction %2 None %79
%270 = OpLabel
%286 = OpVariable %287 Function %288
%273 = OpLoad %15 %47
%274 = OpLoad %16 %49
%275 = OpLoad %18 %54
%276 = OpLoad %20 %58
%277 = OpLoad %24 %64
OpBranch %289
%289 = OpLabel
%290 = OpCompositeExtract %7 %280 0
%292 = OpSampledImage %291 %273 %277
%293 = OpImageSampleImplicitLod %23 %292 %290
%294 = OpLoad %23 %286
%295 = OpFAdd %23 %294 %293
OpStore %286 %295
%297 = OpSampledImage %296 %274 %277
%298 = OpImageSampleImplicitLod %23 %297 %280
%299 = OpLoad %23 %286
%300 = OpFAdd %23 %299 %298
OpStore %286 %300
%301 = OpSampledImage %296 %274 %277
%302 = OpImageSampleImplicitLod %23 %301 %280 ConstOffset %30
%303 = OpLoad %23 %286
%304 = OpFAdd %23 %303 %302
OpStore %286 %304
%305 = OpSampledImage %296 %274 %277
%306 = OpImageSampleExplicitLod %23 %305 %280 Lod %283
%307 = OpLoad %23 %286
%308 = OpFAdd %23 %307 %306
OpStore %286 %308
%309 = OpSampledImage %296 %274 %277
%310 = OpImageSampleExplicitLod %23 %309 %280 Lod|ConstOffset %283 %30
%311 = OpLoad %23 %286
%312 = OpFAdd %23 %311 %310
OpStore %286 %312
%313 = OpSampledImage %296 %274 %277
%314 = OpImageSampleImplicitLod %23 %313 %280 Bias|ConstOffset %284 %30
%315 = OpLoad %23 %286
%316 = OpFAdd %23 %315 %314
OpStore %286 %316
%318 = OpConvertUToF %7 %198
%319 = OpCompositeConstruct %281 %280 %318
%320 = OpSampledImage %317 %275 %277
%321 = OpImageSampleImplicitLod %23 %320 %319
%322 = OpLoad %23 %286
%323 = OpFAdd %23 %322 %321
OpStore %286 %323
%324 = OpConvertUToF %7 %198
%325 = OpCompositeConstruct %281 %280 %324
%326 = OpSampledImage %317 %275 %277
%327 = OpImageSampleImplicitLod %23 %326 %325 ConstOffset %30
%328 = OpLoad %23 %286
%329 = OpFAdd %23 %328 %327
OpStore %286 %329
%330 = OpConvertUToF %7 %198
%331 = OpCompositeConstruct %281 %280 %330
%332 = OpSampledImage %317 %275 %277
%333 = OpImageSampleExplicitLod %23 %332 %331 Lod %283
%334 = OpLoad %23 %286
%335 = OpFAdd %23 %334 %333
OpStore %286 %335
%336 = OpConvertUToF %7 %198
%337 = OpCompositeConstruct %281 %280 %336
%338 = OpSampledImage %317 %275 %277
%339 = OpImageSampleExplicitLod %23 %338 %337 Lod|ConstOffset %283 %30
%340 = OpLoad %23 %286
%341 = OpFAdd %23 %340 %339
OpStore %286 %341
%342 = OpConvertUToF %7 %198
%343 = OpCompositeConstruct %281 %280 %342
%344 = OpSampledImage %317 %275 %277
%345 = OpImageSampleImplicitLod %23 %344 %343 Bias|ConstOffset %284 %30
%346 = OpLoad %23 %286
%347 = OpFAdd %23 %346 %345
OpStore %286 %347
%348 = OpConvertSToF %7 %285
%349 = OpCompositeConstruct %281 %280 %348
%350 = OpSampledImage %317 %275 %277
%351 = OpImageSampleImplicitLod %23 %350 %349
%352 = OpLoad %23 %286
%353 = OpFAdd %23 %352 %351
OpStore %286 %353
%354 = OpConvertSToF %7 %285
%355 = OpCompositeConstruct %281 %280 %354
%356 = OpSampledImage %317 %275 %277
%357 = OpImageSampleImplicitLod %23 %356 %355 ConstOffset %30
%358 = OpLoad %23 %286
%359 = OpFAdd %23 %358 %357
OpStore %286 %359
%360 = OpConvertSToF %7 %285
%361 = OpCompositeConstruct %281 %280 %360
%362 = OpSampledImage %317 %275 %277
%363 = OpImageSampleExplicitLod %23 %362 %361 Lod %283
%364 = OpLoad %23 %286
%365 = OpFAdd %23 %364 %363
OpStore %286 %365
%366 = OpConvertSToF %7 %285
%367 = OpCompositeConstruct %281 %280 %366
%368 = OpSampledImage %317 %275 %277
%369 = OpImageSampleExplicitLod %23 %368 %367 Lod|ConstOffset %283 %30
%370 = OpLoad %23 %286
%371 = OpFAdd %23 %370 %369
OpStore %286 %371
%372 = OpConvertSToF %7 %285
%373 = OpCompositeConstruct %281 %280 %372
%374 = OpSampledImage %317 %275 %277
%375 = OpImageSampleImplicitLod %23 %374 %373 Bias|ConstOffset %284 %30
%376 = OpLoad %23 %286
%377 = OpFAdd %23 %376 %375
OpStore %286 %377
%379 = OpConvertUToF %7 %198
%380 = OpCompositeConstruct %23 %282 %379
%381 = OpSampledImage %378 %276 %277
%382 = OpImageSampleImplicitLod %23 %381 %380
%383 = OpLoad %23 %286
%384 = OpFAdd %23 %383 %382
OpStore %286 %384
%385 = OpConvertUToF %7 %198
%386 = OpCompositeConstruct %23 %282 %385
%387 = OpSampledImage %378 %276 %277
%388 = OpImageSampleExplicitLod %23 %387 %386 Lod %283
%389 = OpLoad %23 %286
%390 = OpFAdd %23 %389 %388
OpStore %286 %390
%391 = OpConvertUToF %7 %198
%392 = OpCompositeConstruct %23 %282 %391
%393 = OpSampledImage %378 %276 %277
%394 = OpImageSampleImplicitLod %23 %393 %392 Bias %284
%395 = OpLoad %23 %286
%396 = OpFAdd %23 %395 %394
OpStore %286 %396
%397 = OpConvertSToF %7 %285
%398 = OpCompositeConstruct %23 %282 %397
%399 = OpSampledImage %378 %276 %277
%400 = OpImageSampleImplicitLod %23 %399 %398
%401 = OpLoad %23 %286
%402 = OpFAdd %23 %401 %400
OpStore %286 %402
%403 = OpConvertSToF %7 %285
%404 = OpCompositeConstruct %23 %282 %403
%405 = OpSampledImage %378 %276 %277
%406 = OpImageSampleExplicitLod %23 %405 %404 Lod %283
%407 = OpLoad %23 %286
%408 = OpFAdd %23 %407 %406
OpStore %286 %408
%409 = OpConvertSToF %7 %285
%410 = OpCompositeConstruct %23 %282 %409
%411 = OpSampledImage %378 %276 %277
%412 = OpImageSampleImplicitLod %23 %411 %410 Bias %284
%413 = OpLoad %23 %286
%414 = OpFAdd %23 %413 %412
OpStore %286 %414
%415 = OpLoad %23 %286
OpStore %271 %415
OpReturn
OpFunctionEnd
%417 = OpFunction %2 None %79
%414 = OpLabel
%422 = OpVariable %423 Function %424
%418 = OpLoad %24 %66
%419 = OpLoad %25 %68
%420 = OpLoad %26 %70
%421 = OpLoad %27 %72
OpBranch %425
%425 = OpLabel
%427 = OpSampledImage %426 %419 %418
%428 = OpImageSampleDrefImplicitLod %7 %427 %278 %276
%429 = OpLoad %7 %422
%430 = OpFAdd %7 %429 %428
OpStore %422 %430
%432 = OpConvertUToF %7 %198
%433 = OpCompositeConstruct %279 %278 %432
%434 = OpSampledImage %431 %420 %418
%435 = OpImageSampleDrefImplicitLod %7 %434 %433 %276
%436 = OpLoad %7 %422
%437 = OpFAdd %7 %436 %435
OpStore %422 %437
%438 = OpConvertSToF %7 %283
%439 = OpCompositeConstruct %279 %278 %438
%440 = OpSampledImage %431 %420 %418
%441 = OpImageSampleDrefImplicitLod %7 %440 %439 %276
%442 = OpLoad %7 %422
%443 = OpFAdd %7 %442 %441
OpStore %422 %443
%445 = OpSampledImage %444 %421 %418
%446 = OpImageSampleDrefImplicitLod %7 %445 %280 %276
%447 = OpLoad %7 %422
%448 = OpFAdd %7 %447 %446
OpStore %422 %448
%449 = OpSampledImage %426 %419 %418
%450 = OpImageSampleDrefExplicitLod %7 %449 %278 %276 Lod %451
%452 = OpLoad %7 %422
%453 = OpFAdd %7 %452 %450
OpStore %422 %453
%454 = OpConvertUToF %7 %198
%455 = OpCompositeConstruct %279 %278 %454
%456 = OpSampledImage %431 %420 %418
%457 = OpImageSampleDrefExplicitLod %7 %456 %455 %276 Lod %451
%458 = OpLoad %7 %422
%459 = OpFAdd %7 %458 %457
OpStore %422 %459
%460 = OpConvertSToF %7 %283
%461 = OpCompositeConstruct %279 %278 %460
%462 = OpSampledImage %431 %420 %418
%463 = OpImageSampleDrefExplicitLod %7 %462 %461 %276 Lod %451
%464 = OpLoad %7 %422
%465 = OpFAdd %7 %464 %463
OpStore %422 %465
%466 = OpSampledImage %444 %421 %418
%467 = OpImageSampleDrefExplicitLod %7 %466 %280 %276 Lod %451
%468 = OpLoad %7 %422
%469 = OpFAdd %7 %468 %467
OpStore %422 %469
%470 = OpLoad %7 %422
OpStore %415 %470
%419 = OpFunction %2 None %79
%416 = OpLabel
%424 = OpVariable %425 Function %426
%420 = OpLoad %24 %66
%421 = OpLoad %25 %68
%422 = OpLoad %26 %70
%423 = OpLoad %27 %72
OpBranch %427
%427 = OpLabel
%429 = OpSampledImage %428 %421 %420
%430 = OpImageSampleDrefImplicitLod %7 %429 %280 %278
%431 = OpLoad %7 %424
%432 = OpFAdd %7 %431 %430
OpStore %424 %432
%434 = OpConvertUToF %7 %198
%435 = OpCompositeConstruct %281 %280 %434
%436 = OpSampledImage %433 %422 %420
%437 = OpImageSampleDrefImplicitLod %7 %436 %435 %278
%438 = OpLoad %7 %424
%439 = OpFAdd %7 %438 %437
OpStore %424 %439
%440 = OpConvertSToF %7 %285
%441 = OpCompositeConstruct %281 %280 %440
%442 = OpSampledImage %433 %422 %420
%443 = OpImageSampleDrefImplicitLod %7 %442 %441 %278
%444 = OpLoad %7 %424
%445 = OpFAdd %7 %444 %443
OpStore %424 %445
%447 = OpSampledImage %446 %423 %420
%448 = OpImageSampleDrefImplicitLod %7 %447 %282 %278
%449 = OpLoad %7 %424
%450 = OpFAdd %7 %449 %448
OpStore %424 %450
%451 = OpSampledImage %428 %421 %420
%452 = OpImageSampleDrefExplicitLod %7 %451 %280 %278 Lod %453
%454 = OpLoad %7 %424
%455 = OpFAdd %7 %454 %452
OpStore %424 %455
%456 = OpConvertUToF %7 %198
%457 = OpCompositeConstruct %281 %280 %456
%458 = OpSampledImage %433 %422 %420
%459 = OpImageSampleDrefExplicitLod %7 %458 %457 %278 Lod %453
%460 = OpLoad %7 %424
%461 = OpFAdd %7 %460 %459
OpStore %424 %461
%462 = OpConvertSToF %7 %285
%463 = OpCompositeConstruct %281 %280 %462
%464 = OpSampledImage %433 %422 %420
%465 = OpImageSampleDrefExplicitLod %7 %464 %463 %278 Lod %453
%466 = OpLoad %7 %424
%467 = OpFAdd %7 %466 %465
OpStore %424 %467
%468 = OpSampledImage %446 %423 %420
%469 = OpImageSampleDrefExplicitLod %7 %468 %282 %278 Lod %453
%470 = OpLoad %7 %424
%471 = OpFAdd %7 %470 %469
OpStore %424 %471
%472 = OpLoad %7 %424
OpStore %417 %472
OpReturn
OpFunctionEnd
%473 = OpFunction %2 None %79
%471 = OpLabel
%474 = OpLoad %16 %49
%475 = OpLoad %3 %51
%476 = OpLoad %17 %52
%477 = OpLoad %24 %64
%478 = OpLoad %24 %66
%479 = OpLoad %25 %68
OpBranch %480
%480 = OpLabel
%481 = OpSampledImage %294 %474 %477
%482 = OpImageGather %23 %481 %278 %483
%484 = OpSampledImage %294 %474 %477
%485 = OpImageGather %23 %484 %278 %486 ConstOffset %30
%487 = OpSampledImage %426 %479 %478
%488 = OpImageDrefGather %23 %487 %278 %276
%489 = OpSampledImage %426 %479 %478
%490 = OpImageDrefGather %23 %489 %278 %276 ConstOffset %30
%492 = OpSampledImage %491 %475 %477
%493 = OpImageGather %98 %492 %278 %198
%496 = OpSampledImage %495 %476 %477
%497 = OpImageGather %494 %496 %278 %198
%498 = OpConvertUToF %23 %493
%499 = OpConvertSToF %23 %497
%500 = OpFAdd %23 %498 %499
%501 = OpFAdd %23 %482 %485
%502 = OpFAdd %23 %501 %488
%503 = OpFAdd %23 %502 %490
%504 = OpFAdd %23 %503 %500
OpStore %472 %504
%475 = OpFunction %2 None %79
%473 = OpLabel
%476 = OpLoad %16 %49
%477 = OpLoad %3 %51
%478 = OpLoad %17 %52
%479 = OpLoad %24 %64
%480 = OpLoad %24 %66
%481 = OpLoad %25 %68
OpBranch %482
%482 = OpLabel
%483 = OpSampledImage %296 %476 %479
%484 = OpImageGather %23 %483 %280 %485
%486 = OpSampledImage %296 %476 %479
%487 = OpImageGather %23 %486 %280 %488 ConstOffset %30
%489 = OpSampledImage %428 %481 %480
%490 = OpImageDrefGather %23 %489 %280 %278
%491 = OpSampledImage %428 %481 %480
%492 = OpImageDrefGather %23 %491 %280 %278 ConstOffset %30
%494 = OpSampledImage %493 %477 %479
%495 = OpImageGather %98 %494 %280 %198
%498 = OpSampledImage %497 %478 %479
%499 = OpImageGather %496 %498 %280 %198
%500 = OpConvertUToF %23 %495
%501 = OpConvertSToF %23 %499
%502 = OpFAdd %23 %500 %501
%503 = OpFAdd %23 %484 %487
%504 = OpFAdd %23 %503 %490
%505 = OpFAdd %23 %504 %492
%506 = OpFAdd %23 %505 %502
OpStore %474 %506
OpReturn
OpFunctionEnd
%507 = OpFunction %2 None %79
%505 = OpLabel
%508 = OpLoad %24 %64
%509 = OpLoad %25 %68
OpBranch %510
%510 = OpLabel
%511 = OpSampledImage %426 %509 %508
%512 = OpImageSampleImplicitLod %23 %511 %278
%513 = OpCompositeExtract %7 %512 0
%514 = OpSampledImage %426 %509 %508
%515 = OpImageGather %23 %514 %278 %198
%516 = OpCompositeConstruct %23 %513 %513 %513 %513
%517 = OpFAdd %23 %516 %515
OpStore %506 %517
%509 = OpFunction %2 None %79
%507 = OpLabel
%510 = OpLoad %24 %64
%511 = OpLoad %25 %68
OpBranch %512
%512 = OpLabel
%513 = OpSampledImage %428 %511 %510
%514 = OpImageSampleImplicitLod %23 %513 %280
%515 = OpCompositeExtract %7 %514 0
%516 = OpSampledImage %428 %511 %510
%517 = OpImageGather %23 %516 %280 %198
%518 = OpCompositeConstruct %23 %515 %515 %515 %515
%519 = OpFAdd %23 %518 %517
OpStore %508 %519
OpReturn
OpFunctionEnd

156
naga/tests/out/spv/math-functions.spvasm
View File

@ -1,12 +1,12 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 87
; Bound: 95
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %17 "main"
OpExecutionMode %17 OriginUpperLeft
OpEntryPoint Fragment %18 "main"
OpExecutionMode %18 OriginUpperLeft
OpMemberDecorate %11 0 Offset 0
OpMemberDecorate %11 1 Offset 4
OpMemberDecorate %12 0 Offset 0
@ -31,77 +31,85 @@ OpMemberDecorate %15 1 Offset 16
%13 = OpTypeStruct %4 %4
%14 = OpTypeStruct %3 %6
%15 = OpTypeStruct %4 %5
%18 = OpTypeFunction %2
%19 = OpConstant %3 1.0
%20 = OpConstant %3 0.0
%21 = OpConstantComposite %4 %20 %20 %20 %20
%22 = OpConstant %6 -1
%23 = OpConstantComposite %5 %22 %22 %22 %22
%24 = OpConstant %3 -1.0
%25 = OpConstantComposite %4 %24 %24 %24 %24
%26 = OpConstantNull %7
%27 = OpConstant %9 4294967295
%28 = OpConstantComposite %7 %22 %22
%29 = OpConstant %9 0
%30 = OpConstantComposite %8 %29 %29
%31 = OpConstant %6 0
%32 = OpConstantComposite %7 %31 %31
%33 = OpConstant %9 32
%34 = OpConstant %6 32
%35 = OpConstantComposite %8 %33 %33
%36 = OpConstantComposite %7 %34 %34
%37 = OpConstant %9 31
%38 = OpConstantComposite %8 %37 %37
%39 = OpConstant %6 2
%40 = OpConstant %3 2.0
%41 = OpConstantComposite %10 %19 %40
%42 = OpConstant %6 3
%43 = OpConstant %6 4
%44 = OpConstantComposite %7 %42 %43
%45 = OpConstant %3 1.5
%46 = OpConstantComposite %10 %45 %45
%47 = OpConstantComposite %4 %45 %45 %45 %45
%54 = OpConstantComposite %4 %19 %19 %19 %19
%57 = OpConstantNull %6
%17 = OpFunction %2 None %18
%16 = OpLabel
OpBranch %48
%48 = OpLabel
%49 = OpExtInst %3 %1 Degrees %19
%50 = OpExtInst %3 %1 Radians %19
%51 = OpExtInst %4 %1 Degrees %21
%52 = OpExtInst %4 %1 Radians %21
%53 = OpExtInst %4 %1 FClamp %21 %21 %54
%55 = OpExtInst %4 %1 Refract %21 %21 %19
%58 = OpCompositeExtract %6 %26 0
%59 = OpCompositeExtract %6 %26 0
%60 = OpIMul %6 %58 %59
%61 = OpIAdd %6 %57 %60
%62 = OpCompositeExtract %6 %26 1
%63 = OpCompositeExtract %6 %26 1
%16 = OpTypeVector %3 3
%19 = OpTypeFunction %2
%20 = OpConstant %3 1.0
%21 = OpConstant %3 0.0
%22 = OpConstantComposite %4 %21 %21 %21 %21
%23 = OpConstant %6 -1
%24 = OpConstantComposite %5 %23 %23 %23 %23
%25 = OpConstant %3 -1.0
%26 = OpConstantComposite %4 %25 %25 %25 %25
%27 = OpConstantNull %7
%28 = OpConstant %9 4294967295
%29 = OpConstantComposite %7 %23 %23
%30 = OpConstant %9 0
%31 = OpConstantComposite %8 %30 %30
%32 = OpConstant %6 0
%33 = OpConstantComposite %7 %32 %32
%34 = OpConstant %9 32
%35 = OpConstant %6 32
%36 = OpConstantComposite %8 %34 %34
%37 = OpConstantComposite %7 %35 %35
%38 = OpConstant %9 31
%39 = OpConstantComposite %8 %38 %38
%40 = OpConstant %6 2
%41 = OpConstant %3 2.0
%42 = OpConstantComposite %10 %20 %41
%43 = OpConstant %6 3
%44 = OpConstant %6 4
%45 = OpConstantComposite %7 %43 %44
%46 = OpConstant %3 1.5
%47 = OpConstantComposite %10 %46 %46
%48 = OpConstantComposite %4 %46 %46 %46 %46
%49 = OpConstantComposite %10 %20 %20
%50 = OpConstantComposite %16 %20 %20 %20
%51 = OpConstantComposite %4 %20 %20 %20 %20
%58 = OpConstantComposite %4 %20 %20 %20 %20
%61 = OpConstantNull %6
%18 = OpFunction %2 None %19
%17 = OpLabel
OpBranch %52
%52 = OpLabel
%53 = OpExtInst %3 %1 Degrees %20
%54 = OpExtInst %3 %1 Radians %20
%55 = OpExtInst %4 %1 Degrees %22
%56 = OpExtInst %4 %1 Radians %22
%57 = OpExtInst %4 %1 FClamp %22 %22 %58
%59 = OpExtInst %4 %1 Refract %22 %22 %20
%62 = OpCompositeExtract %6 %27 0
%63 = OpCompositeExtract %6 %27 0
%64 = OpIMul %6 %62 %63
%56 = OpIAdd %6 %61 %64
%65 = OpExtInst %3 %1 Ldexp %19 %39
%66 = OpExtInst %10 %1 Ldexp %41 %44
%67 = OpExtInst %11 %1 ModfStruct %45
%68 = OpExtInst %11 %1 ModfStruct %45
%69 = OpCompositeExtract %3 %68 0
%70 = OpExtInst %11 %1 ModfStruct %45
%71 = OpCompositeExtract %3 %70 1
%72 = OpExtInst %12 %1 ModfStruct %46
%73 = OpExtInst %13 %1 ModfStruct %47
%74 = OpCompositeExtract %4 %73 1
%75 = OpCompositeExtract %3 %74 0
%76 = OpExtInst %12 %1 ModfStruct %46
%77 = OpCompositeExtract %10 %76 0
%78 = OpCompositeExtract %3 %77 1
%79 = OpExtInst %14 %1 FrexpStruct %45
%80 = OpExtInst %14 %1 FrexpStruct %45
%81 = OpCompositeExtract %3 %80 0
%82 = OpExtInst %14 %1 FrexpStruct %45
%83 = OpCompositeExtract %6 %82 1
%84 = OpExtInst %15 %1 FrexpStruct %47
%85 = OpCompositeExtract %5 %84 1
%86 = OpCompositeExtract %6 %85 0
%65 = OpIAdd %6 %61 %64
%66 = OpCompositeExtract %6 %27 1
%67 = OpCompositeExtract %6 %27 1
%68 = OpIMul %6 %66 %67
%60 = OpIAdd %6 %65 %68
%69 = OpExtInst %3 %1 Ldexp %20 %40
%70 = OpExtInst %10 %1 Ldexp %42 %45
%71 = OpExtInst %11 %1 ModfStruct %46
%72 = OpExtInst %11 %1 ModfStruct %46
%73 = OpCompositeExtract %3 %72 0
%74 = OpExtInst %11 %1 ModfStruct %46
%75 = OpCompositeExtract %3 %74 1
%76 = OpExtInst %12 %1 ModfStruct %47
%77 = OpExtInst %13 %1 ModfStruct %48
%78 = OpCompositeExtract %4 %77 1
%79 = OpCompositeExtract %3 %78 0
%80 = OpExtInst %12 %1 ModfStruct %47
%81 = OpCompositeExtract %10 %80 0
%82 = OpCompositeExtract %3 %81 1
%83 = OpExtInst %14 %1 FrexpStruct %46
%84 = OpExtInst %14 %1 FrexpStruct %46
%85 = OpCompositeExtract %3 %84 0
%86 = OpExtInst %14 %1 FrexpStruct %46
%87 = OpCompositeExtract %6 %86 1
%88 = OpExtInst %15 %1 FrexpStruct %48
%89 = OpCompositeExtract %5 %88 1
%90 = OpCompositeExtract %6 %89 0
%91 = OpQuantizeToF16 %3 %20
%92 = OpQuantizeToF16 %10 %49
%93 = OpQuantizeToF16 %16 %50
%94 = OpQuantizeToF16 %4 %51
OpReturn
OpFunctionEnd

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