nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-23 07:14:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Ray Queries (#6291)

Browse Source 
* fix more compile errors

* fix err

* fix all compile errors

* fix clippy errors

* fix issues in ray-traced-triangle example

* fix clippy warnings

* fix wasm not compiling

* fix test

* fix clippy (again)

* fix some rustfmt issues

* fix more rustfmt issues

* fix even more rustfmt issues

* fix yet more rustfmt issues

* add #[allow(dead_code)] to some structs to fix clippy on web assembly

* more format fixes

* even more format fixes

* close to final format fixes

* final format fixes (hopefully)

* actual final format fixes (hopefully)

* move ray_cube_cube example and fix bugs

* move other examples

* fix crash when dropping

* fix test

* apply clippy and rustfmt changes

* fix some memory leaks

* fix incompatibility with new wgpu trunk changes

* fix clippy & fmt

* fix clippy errors on github

* update to trunk

* remove a line that printed log "Created blas ..." in device create tlas

* fix merge issues

* fix merge

* fix fmt

* restored example ray scene functionality

* WIP fix rebase

* Fix compile issues

* More fixes

* Fix rebase

* Add wgls backend write_value_type AccelerationStructure

* Add back xtask Cargo.lock

* Revert serde path

* Doc typos

* Add -diff

* Revert more serde

* Revert serde again?

* Fix doc typo

* add test

* clippy & fmt

* fix error "blas does not exist"

* fix various things

* fix unused acceleration structures causing invalid SPIR-V

* fix merge

* fix use after free in a different way (now moved tlas instances to wgpu_core)

* refactor & fmt

* remove accidentally created file

* remove accidentally created file (actually)

* merge & fix

* Fix problems

* fix validation error

* fix validation

* fix merge

* fmt

* fix merge

* use staging buffer trackers in creating the ray-tracing staging buffers (not the tlas_s)

* change one missed staging buffer's tracker

* wip: example

* finished shadow!

* fix merge

* clippy

* remove some uses of blas ids

* fmt

* fix merge

* deduplicate code shared between command_encoder_build_acceleration_structures and command_encoder_build_acceleration_structures_unsafe_tlas

* clippy

* fix merge

* fix warnings

* clippy & fmt

* re-add matching to get_raw_tlas_instance_size and tlas_instance_into_bytes

* fmt

* fix command

* wip: fix merge

* fix merge

* fmt & docs

* fix wasm clippy

* fix issues with tests

* fix issues with tests

* fmt & clippy (wasm)

* wip ray-traced-triangle (non hal)

* unify camera location in ray_shadows

* fix tests

* add readme to ray_traced_triangle

* note that vertex format must be f32x3 (not yet validated)

* changelog

* require VertexFormat::Float32x3

* fix compile error

* when creating and building acceleration structures require Features::RAY_TRACING_ACCELERATION_STRUCTURE is enabled

* remove duplicate workaround

* fix most feedback

* remove all changes from .gitattributes

* fix compile error in test

* more changes along same lines

* format

* apply feedback to more places

* fix `cargo clippy` not just `cargo clippy --all-features`

* cache max instances from tlas creation

* fix clippy

* fix examples

* fix tests

* remove now useless docs

* add more docs

* yet more docs

* still more docs

* Documentation.

* Running out of unique commit messages to say writing docs.

* Add more references to docs.

* Get BindGroup in scope

* Remove unintentional changes

* Remove #[allow(dead_code)]

* Remove old comments

* Add device.poll

* make ray_scene start right way up

* fmt

* move scratch buffer to its own module

* actually remove intentional changes (turns out wgt is inconsistent)

* use better names in ray_scene/shader.wgsl

* fmt

* remove more unintentional changes.

* work on ray_tracing use after free

* spelling

Co-authored-by: Connor Fitzgerald <connorwadefitzgerald@gmail.com>

* grammar

Co-authored-by: Connor Fitzgerald <connorwadefitzgerald@gmail.com>

* move things into struct (should be all vectors with 3 item types or more)

* remove long-standing locks.

* fmt

* move not using index buffer to test

* impl index & index mut for tlas package

* impl index & index mut for tlas package

* fix errors

* fix test errors

* move everything out of ray_tracing.rs

* more index_mut

* fmt

* fix docs links

* move to separate impl block

* fix older request review

* fmt

* remove missed deref

* fix wasm clippy

* change to suggestions

* fmt

* add a note as to why `command_encoder_build_acceleration_structures_unsafe_tlas` and `command_encoder_build_acceleration_structures` are currently separate

* add a todo to reconsider soon `command_encoder_build_acceleration_structures_unsafe_tlas`

* switch to mat4 in ray-tracing examples

* Combine RT Tests Under Single Submodule

* Add More Raytracing Tests and Re-Organize

* Documentation and Formatting Fixes

* Remove Forced Alignment in wgpu-hal

* work on out_of_order_as_build and add out_of_order_as_build_use

* fmt

* update test label naming

* fix one validation error

* fix clippy & fmt

* add acceleration structure creation tests

* make mis-matched index test use the correct function

* move validation into validate_command_buffer

* fix tests

* add blas and tlas to pending writes to remove (now unneeded) `device.poll()`s

* remove commented out code from examples and properly use the constants in ray_shadow

* add examples to examples\README.md

* rename ray tracing to ray queries in readme

* remove unneeded usage from scratch buffer creation

* use snatchables for raw acceleration structures

* format

* add EXPERIMENTAL_ to features

* fix missed comment

---------

Co-authored-by: Daniel Keitel <daniel.mg.keitel@fau.de>
Co-authored-by: JMS55 <47158642+JMS55@users.noreply.github.com>
Co-authored-by: atlas dostal <rodol@rivalrebels.com>
Co-authored-by: vero <email@atlasdostal.com>
Co-authored-by: Connor Fitzgerald <connorwadefitzgerald@gmail.com>
This commit is contained in:
Vecvec 2024-11-10 06:59:48 +13:00 committed by GitHub
parent 01973a98dc
commit 79a6f2cd31
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
77 changed files with 10364 additions and 81 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
CHANGELOG.md
View File

@ -200,6 +200,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).

4
Cargo.lock generated
View File

@ -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"
@ -3746,6 +3749,7 @@ dependencies = [
"ctor",
"env_logger",
"futures-lite",
"glam",
"image",
"itertools",
"js-sys",

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);
}

106
player/src/lib.rs
View File

@ -119,6 +119,94 @@ impl GlobalPlay for wgc::global::Global {
)
.unwrap();
}
trace::Command::BuildAccelerationStructuresUnsafeTlas { blas, tlas } => {
let blas_iter = blas.iter().map(|x| {
let geometries = match &x.geometries {
wgc::ray_tracing::TraceBlasGeometries::TriangleGeometries(
triangle_geometries,
) => {
let iter = triangle_geometries.iter().map(|tg| {
wgc::ray_tracing::BlasTriangleGeometry {
size: &tg.size,
vertex_buffer: tg.vertex_buffer,
index_buffer: tg.index_buffer,
transform_buffer: tg.transform_buffer,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
transform_buffer_offset: tg.transform_buffer_offset,
}
});
wgc::ray_tracing::BlasGeometries::TriangleGeometries(Box::new(iter))
}
};
wgc::ray_tracing::BlasBuildEntry {
blas_id: x.blas_id,
geometries,
}
});
if !tlas.is_empty() {
log::error!("a trace of command_encoder_build_acceleration_structures_unsafe_tlas containing a tlas build is not replayable! skipping tlas build");
}
self.command_encoder_build_acceleration_structures_unsafe_tlas(
encoder,
blas_iter,
std::iter::empty(),
)
.unwrap();
}
trace::Command::BuildAccelerationStructures { blas, tlas } => {
let blas_iter = blas.iter().map(|x| {
let geometries = match &x.geometries {
wgc::ray_tracing::TraceBlasGeometries::TriangleGeometries(
triangle_geometries,
) => {
let iter = triangle_geometries.iter().map(|tg| {
wgc::ray_tracing::BlasTriangleGeometry {
size: &tg.size,
vertex_buffer: tg.vertex_buffer,
index_buffer: tg.index_buffer,
transform_buffer: tg.transform_buffer,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
transform_buffer_offset: tg.transform_buffer_offset,
}
});
wgc::ray_tracing::BlasGeometries::TriangleGeometries(Box::new(iter))
}
};
wgc::ray_tracing::BlasBuildEntry {
blas_id: x.blas_id,
geometries,
}
});
let tlas_iter = tlas.iter().map(|x| {
let instances = x.instances.iter().map(|instance| {
instance
.as_ref()
.map(|instance| wgc::ray_tracing::TlasInstance {
blas_id: instance.blas_id,
transform: &instance.transform,
custom_index: instance.custom_index,
mask: instance.mask,
})
});
wgc::ray_tracing::TlasPackage {
tlas_id: x.tlas_id,
instances: Box::new(instances),
lowest_unmodified: x.lowest_unmodified,
}
});
self.command_encoder_build_acceleration_structures(
encoder, blas_iter, tlas_iter,
)
.unwrap();
}
}
}
let (cmd_buf, error) =
@ -360,6 +448,24 @@ impl GlobalPlay for wgc::global::Global {
let cmdbuf = self.encode_commands(encoder, commands);
self.queue_submit(queue, &[cmdbuf]).unwrap();
}
Action::CreateBlas { id, desc, sizes } => {
self.device_create_blas(device, &desc, sizes, Some(id));
}
Action::FreeBlas(id) => {
self.blas_destroy(id).unwrap();
}
Action::DestroyBlas(id) => {
self.blas_drop(id);
}
Action::CreateTlas { id, desc } => {
self.device_create_tlas(device, &desc, Some(id));
}
Action::FreeTlas(id) => {
self.tlas_destroy(id).unwrap();
}
Action::DestroyTlas(id) => {
self.tlas_drop(id);
}
}
}
}

1
tests/Cargo.toml
View File

@ -40,6 +40,7 @@ strum = { workspace = true, features = ["derive"] }
wgpu-macros.workspace = true
wgpu = { workspace = true, features = ["wgsl"] }
wgt = { workspace = true, features = ["serde"] }
glam.workspace = true
[target.'cfg(not(target_arch = "wasm32"))'.dependencies]
env_logger.workspace = true

3
tests/tests/buffer_usages.rs
View File

@ -21,7 +21,8 @@ const NEEDS_MAPPABLE_PRIMARY_BUFFERS: &[Bu; 7] = &[
Bu::MAP_WRITE.union(Bu::MAP_READ),
Bu::MAP_READ.union(Bu::COPY_DST.union(Bu::STORAGE)),
Bu::MAP_WRITE.union(Bu::COPY_SRC.union(Bu::STORAGE)),
Bu::all(),
// these two require acceleration_structures feature
Bu::all().intersection(Bu::BLAS_INPUT.union(Bu::TLAS_INPUT).complement()),
];
const INVALID_BITS: Bu = Bu::from_bits_retain(0b1111111111111);
const ALWAYS_FAIL: &[Bu; 2] = &[Bu::empty(), INVALID_BITS];

285
tests/tests/ray_tracing/as_build.rs Normal file
View File

@ -0,0 +1,285 @@
use std::mem;
use wgpu::{
util::{BufferInitDescriptor, DeviceExt},
*,
};
use wgpu_test::{fail, gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
struct AsBuildContext {
vertices: Buffer,
blas_size: BlasTriangleGeometrySizeDescriptor,
blas: Blas,
// Putting this last, forces the BLAS to die before the TLAS.
tlas_package: TlasPackage,
}
impl AsBuildContext {
fn new(ctx: &TestingContext) -> Self {
let vertices = ctx.device.create_buffer_init(&BufferInitDescriptor {
label: None,
contents: &[0; mem::size_of::<[[f32; 3]; 3]>()],
usage: BufferUsages::BLAS_INPUT,
});
let blas_size = BlasTriangleGeometrySizeDescriptor {
vertex_format: VertexFormat::Float32x3,
vertex_count: 3,
index_format: None,
index_count: None,
flags: AccelerationStructureGeometryFlags::empty(),
};
let blas = ctx.device.create_blas(
&CreateBlasDescriptor {
label: Some("BLAS"),
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size.clone()],
},
);
let tlas = ctx.device.create_tlas(&CreateTlasDescriptor {
label: Some("TLAS"),
max_instances: 1,
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
});
let mut tlas_package = TlasPackage::new(tlas);
tlas_package[0] = Some(TlasInstance::new(
&blas,
[1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0],
0,
0xFF,
));
Self {
vertices,
blas_size,
blas,
tlas_package,
}
}
fn blas_build_entry(&self) -> BlasBuildEntry {
BlasBuildEntry {
blas: &self.blas,
geometry: BlasGeometries::TriangleGeometries(vec![BlasTriangleGeometry {
size: &self.blas_size,
vertex_buffer: &self.vertices,
first_vertex: 0,
vertex_stride: mem::size_of::<[f32; 3]>() as BufferAddress,
index_buffer: None,
index_buffer_offset: None,
transform_buffer: None,
transform_buffer_offset: None,
}]),
}
}
}
#[gpu_test]
static UNBUILT_BLAS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(unbuilt_blas);
fn unbuilt_blas(ctx: TestingContext) {
let as_ctx = AsBuildContext::new(&ctx);
// Build the TLAS package with an unbuilt BLAS.
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
encoder.build_acceleration_structures([], [&as_ctx.tlas_package]);
fail(
&ctx.device,
|| {
ctx.queue.submit([encoder.finish()]);
},
None,
);
}
#[gpu_test]
static OUT_OF_ORDER_AS_BUILD: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(out_of_order_as_build);
fn out_of_order_as_build(ctx: TestingContext) {
let as_ctx = AsBuildContext::new(&ctx);
//
// Encode the TLAS build before the BLAS build, but submit them in the right order.
//
let mut encoder_tlas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("TLAS 1"),
});
encoder_tlas.build_acceleration_structures([], [&as_ctx.tlas_package]);
let mut encoder_blas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("BLAS 1"),
});
encoder_blas.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
ctx.queue
.submit([encoder_blas.finish(), encoder_tlas.finish()]);
drop(as_ctx);
//
// Create a clean `AsBuildContext`
//
let as_ctx = AsBuildContext::new(&ctx);
//
// Encode the BLAS build before the TLAS build, but submit them in the wrong order.
//
let mut encoder_blas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("BLAS 2"),
});
encoder_blas.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
let mut encoder_tlas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("TLAS 2"),
});
encoder_tlas.build_acceleration_structures([], [&as_ctx.tlas_package]);
fail(
&ctx.device,
|| {
ctx.queue
.submit([encoder_tlas.finish(), encoder_blas.finish()]);
},
None,
);
}
#[gpu_test]
static OUT_OF_ORDER_AS_BUILD_USE: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(TestParameters::default().test_features_limits().features(
wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
| wgpu::Features::EXPERIMENTAL_RAY_QUERY,
))
.run_sync(out_of_order_as_build_use);
fn out_of_order_as_build_use(ctx: TestingContext) {
//
// Create a clean `AsBuildContext`
//
let as_ctx = AsBuildContext::new(&ctx);
//
// Build in the right order, then rebuild the BLAS so the TLAS is invalid, then use the TLAS.
//
let mut encoder_blas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("BLAS 1"),
});
encoder_blas.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
let mut encoder_tlas = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("TLAS 1"),
});
encoder_tlas.build_acceleration_structures([], [&as_ctx.tlas_package]);
let mut encoder_blas2 = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("BLAS 2"),
});
encoder_blas2.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
ctx.queue.submit([
encoder_blas.finish(),
encoder_tlas.finish(),
encoder_blas2.finish(),
]);
//
// Create shader to use tlas with
//
let shader = ctx
.device
.create_shader_module(include_wgsl!("shader.wgsl"));
let compute_pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: None,
layout: None,
module: &shader,
entry_point: Some("comp_main"),
compilation_options: Default::default(),
cache: None,
});
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: None,
layout: &compute_pipeline.get_bind_group_layout(0),
entries: &[BindGroupEntry {
binding: 0,
resource: BindingResource::AccelerationStructure(as_ctx.tlas_package.tlas()),
}],
});
//
// Use TLAS
//
let mut encoder_compute = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
{
let mut pass = encoder_compute.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
pass.set_pipeline(&compute_pipeline);
pass.set_bind_group(0, Some(&bind_group), &[]);
pass.dispatch_workgroups(1, 1, 1)
}
fail(
&ctx.device,
|| {
ctx.queue.submit(Some(encoder_compute.finish()));
},
None,
);
}

117
tests/tests/ray_tracing/as_create.rs Normal file
View File

@ -0,0 +1,117 @@
use wgpu::{
AccelerationStructureFlags, AccelerationStructureGeometryFlags,
AccelerationStructureUpdateMode, BlasGeometrySizeDescriptors,
BlasTriangleGeometrySizeDescriptor, CreateBlasDescriptor,
};
use wgpu_macros::gpu_test;
use wgpu_test::{fail, GpuTestConfiguration, TestParameters, TestingContext};
use wgt::{IndexFormat, VertexFormat};
#[gpu_test]
static BLAS_INVALID_VERTEX_FORMAT: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(invalid_vertex_format_blas_create);
fn invalid_vertex_format_blas_create(ctx: TestingContext) {
//
// Create a BLAS with a format that is not allowed
//
let blas_size = BlasTriangleGeometrySizeDescriptor {
vertex_format: VertexFormat::Float32x4,
vertex_count: 3,
index_format: None,
index_count: None,
flags: AccelerationStructureGeometryFlags::empty(),
};
fail(
&ctx.device,
|| {
let _ = ctx.device.create_blas(
&CreateBlasDescriptor {
label: Some("BLAS"),
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size.clone()],
},
);
},
None,
);
}
#[gpu_test]
static BLAS_MISMATCHED_INDEX: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(mismatched_index_blas_create);
fn mismatched_index_blas_create(ctx: TestingContext) {
//
// Create a BLAS with just an index format
//
let blas_size = BlasTriangleGeometrySizeDescriptor {
vertex_format: VertexFormat::Float32x3,
vertex_count: 3,
index_format: Some(IndexFormat::Uint32),
index_count: None,
flags: AccelerationStructureGeometryFlags::empty(),
};
fail(
&ctx.device,
|| {
let _ = ctx.device.create_blas(
&CreateBlasDescriptor {
label: Some("BLAS1"),
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size.clone()],
},
);
},
None,
);
//
// Create a BLAS with just an index count
//
let blas_size = BlasTriangleGeometrySizeDescriptor {
vertex_format: VertexFormat::Float32x3,
vertex_count: 3,
index_format: None,
index_count: Some(3),
flags: AccelerationStructureGeometryFlags::empty(),
};
fail(
&ctx.device,
|| {
let _ = ctx.device.create_blas(
&CreateBlasDescriptor {
label: Some("BLAS2"),
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size.clone()],
},
);
},
None,
);
}

152
tests/tests/ray_tracing/as_use_after_free.rs Normal file
View File

@ -0,0 +1,152 @@
use std::{iter, mem};
use wgpu::{
include_wgsl,
util::{BufferInitDescriptor, DeviceExt},
AccelerationStructureFlags, AccelerationStructureGeometryFlags,
AccelerationStructureUpdateMode, BindGroupDescriptor, BindGroupEntry, BindingResource,
BlasBuildEntry, BlasGeometries, BlasGeometrySizeDescriptors, BlasTriangleGeometry,
BlasTriangleGeometrySizeDescriptor, BufferAddress, BufferUsages, CommandEncoderDescriptor,
ComputePassDescriptor, ComputePipelineDescriptor, CreateBlasDescriptor, CreateTlasDescriptor,
Maintain, TlasInstance, TlasPackage, VertexFormat,
};
use wgpu_macros::gpu_test;
use wgpu_test::{GpuTestConfiguration, TestParameters, TestingContext};
fn required_features() -> wgpu::Features {
wgpu::Features::EXPERIMENTAL_RAY_QUERY
| wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
}
/// This test creates a blas, puts a reference to it in a tlas instance inside a tlas package,
/// drops the blas, and ensures it gets kept alive by the tlas instance. Then it uses the built
/// package in a bindgroup, drops it, and checks that it is kept alive by the bindgroup by
/// executing a shader using that bindgroup.
fn acceleration_structure_use_after_free(ctx: TestingContext) {
// Dummy vertex buffer.
let vertices = ctx.device.create_buffer_init(&BufferInitDescriptor {
label: None,
contents: &[0; mem::size_of::<[[f32; 3]; 3]>()],
usage: BufferUsages::BLAS_INPUT,
});
// Create a BLAS with a single triangle.
let blas_size = BlasTriangleGeometrySizeDescriptor {
vertex_format: VertexFormat::Float32x3,
vertex_count: 3,
index_format: None,
index_count: None,
flags: AccelerationStructureGeometryFlags::empty(),
};
let blas = ctx.device.create_blas(
&CreateBlasDescriptor {
label: Some("blas use after free"),
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
},
BlasGeometrySizeDescriptors::Triangles {
descriptors: vec![blas_size.clone()],
},
);
// Create the TLAS
let tlas = ctx.device.create_tlas(&CreateTlasDescriptor {
label: Some("tlas use after free"),
max_instances: 1,
flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
update_mode: AccelerationStructureUpdateMode::Build,
});
// Put an unbuilt BLAS in the tlas package.
let mut tlas_package = TlasPackage::new(tlas);
tlas_package[0] = Some(TlasInstance::new(
&blas,
[1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0],
0,
0xFF,
));
// Actually build the BLAS.
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
encoder.build_acceleration_structures(
iter::once(&BlasBuildEntry {
blas: &blas,
geometry: BlasGeometries::TriangleGeometries(vec![BlasTriangleGeometry {
size: &blas_size,
vertex_buffer: &vertices,
first_vertex: 0,
vertex_stride: mem::size_of::<[f32; 3]>() as BufferAddress,
index_buffer: None,
index_buffer_offset: None,
transform_buffer: None,
transform_buffer_offset: None,
}]),
}),
iter::empty(),
);
ctx.queue.submit(Some(encoder.finish()));
// Drop the blas and ensure that if it was going to die, it is dead.
drop(blas);
ctx.device.poll(Maintain::Wait);
// build the tlas package to ensure the blas is dropped
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
encoder.build_acceleration_structures(iter::empty(), iter::once(&tlas_package));
ctx.queue.submit(Some(encoder.finish()));
// Create a compute shader that uses an AS.
let shader = ctx
.device
.create_shader_module(include_wgsl!("shader.wgsl"));
let compute_pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: None,
layout: None,
module: &shader,
entry_point: Some("comp_main"),
compilation_options: Default::default(),
cache: None,
});
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: None,
layout: &compute_pipeline.get_bind_group_layout(0),
entries: &[BindGroupEntry {
binding: 0,
resource: BindingResource::AccelerationStructure(tlas_package.tlas()),
}],
});
// Drop the TLAS package and ensure that if it was going to die, it is dead.
drop(tlas_package);
ctx.device.poll(Maintain::Wait);
// Run the pass with the bind group that references the TLAS package.
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
{
let mut pass = encoder.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
pass.set_pipeline(&compute_pipeline);
pass.set_bind_group(0, Some(&bind_group), &[]);
pass.dispatch_workgroups(1, 1, 1)
}
ctx.queue.submit(Some(encoder.finish()));
}
#[gpu_test]
static ACCELERATION_STRUCTURE_USE_AFTER_FREE: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(required_features()),
)
.run_sync(acceleration_structure_use_after_free);

4
tests/tests/ray_tracing/mod.rs Normal file
View File

@ -0,0 +1,4 @@
mod as_build;
mod as_create;
mod as_use_after_free;
mod scene;

83
tests/tests/ray_tracing/scene/mesh_gen.rs Normal file
View File

@ -0,0 +1,83 @@
use bytemuck::{Pod, Zeroable};
use glam::Affine3A;
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
pub 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],
}
}
pub 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())
}
pub 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,
]
}

121
tests/tests/ray_tracing/scene/mod.rs Normal file
View File

@ -0,0 +1,121 @@
use std::{iter, mem};
use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
use wgpu::util::DeviceExt;
use glam::{Affine3A, Quat, Vec3};
mod mesh_gen;
fn acceleration_structure_build(ctx: &TestingContext, use_index_buffer: bool) {
let max_instances = 1000;
let device = &ctx.device;
let (vertex_data, index_data) = mesh_gen::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_buffer = 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 index_format = wgpu::IndexFormat::Uint16;
let index_count = index_data.len() as u32;
let blas_geo_size_desc = wgpu::BlasTriangleGeometrySizeDescriptor {
vertex_format: wgpu::VertexFormat::Float32x3,
vertex_count: vertex_data.len() as u32,
index_format: use_index_buffer.then_some(index_format),
index_count: use_index_buffer.then_some(index_count),
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,
});
let mut tlas_package = wgpu::TlasPackage::new(tlas);
for j in 0..max_instances {
tlas_package[j as usize] = Some(wgpu::TlasInstance::new(
&blas,
mesh_gen::affine_to_rows(&Affine3A::from_rotation_translation(
Quat::from_rotation_y(45.9_f32.to_radians()),
Vec3 {
x: j as f32,
y: 0.0,
z: 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::<mesh_gen::Vertex>() as u64,
index_buffer: use_index_buffer.then_some(&index_buffer),
index_buffer_offset: use_index_buffer.then_some(0),
transform_buffer: None,
transform_buffer_offset: None,
}]),
}),
iter::once(&tlas_package),
);
ctx.queue.submit(Some(encoder.finish()));
ctx.device.poll(wgpu::Maintain::Wait);
}
#[gpu_test]
static ACCELERATION_STRUCTURE_BUILD_NO_INDEX: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(|ctx| {
acceleration_structure_build(&ctx, false);
});
#[gpu_test]
static ACCELERATION_STRUCTURE_BUILD_WITH_INDEX: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.features(wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE),
)
.run_sync(|ctx| {
acceleration_structure_build(&ctx, true);
});

11
tests/tests/ray_tracing/shader.wgsl Normal file
View File

@ -0,0 +1,11 @@
@group(0) @binding(0)
var acc_struct: acceleration_structure;
@workgroup_size(1)
@compute
fn comp_main() {
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}

1
tests/tests/root.rs
View File

@ -37,6 +37,7 @@ mod poll;
mod push_constants;
mod query_set;
mod queue_transfer;
mod ray_tracing;
mod render_pass_ownership;
mod resource_descriptor_accessor;
mod resource_error;

10
wgpu-core/src/binding_model.rs
View File

@ -2,7 +2,7 @@ use crate::{
device::{
bgl, Device, DeviceError, MissingDownlevelFlags, MissingFeatures, SHADER_STAGE_COUNT,
},
id::{BindGroupLayoutId, BufferId, SamplerId, TextureViewId},
id::{BindGroupLayoutId, BufferId, SamplerId, TextureViewId, TlasId},
init_tracker::{BufferInitTrackerAction, TextureInitTrackerAction},
pipeline::{ComputePipeline, RenderPipeline},
resource::{
@ -29,6 +29,7 @@ use std::{
sync::{Arc, OnceLock, Weak},
};
use crate::resource::Tlas;
use thiserror::Error;
#[derive(Clone, Debug, Error)]
@ -302,6 +303,7 @@ pub(crate) struct BindingTypeMaxCountValidator {
storage_buffers: PerStageBindingTypeCounter,
storage_textures: PerStageBindingTypeCounter,
uniform_buffers: PerStageBindingTypeCounter,
acceleration_structures: PerStageBindingTypeCounter,
}
impl BindingTypeMaxCountValidator {
@ -337,7 +339,9 @@ impl BindingTypeMaxCountValidator {
wgt::BindingType::StorageTexture { .. } => {
self.storage_textures.add(binding.visibility, count);
}
wgt::BindingType::AccelerationStructure => todo!(),
wgt::BindingType::AccelerationStructure => {
self.acceleration_structures.add(binding.visibility, count);
}
}
}
@ -781,6 +785,7 @@ pub enum BindingResource<'a> {
SamplerArray(Cow<'a, [SamplerId]>),
TextureView(TextureViewId),
TextureViewArray(Cow<'a, [TextureViewId]>),
AccelerationStructure(TlasId),
}
// Note: Duplicated in `wgpu-rs` as `BindingResource`
@ -793,6 +798,7 @@ pub enum ResolvedBindingResource<'a> {
SamplerArray(Cow<'a, [Arc<Sampler>]>),
TextureView(Arc<TextureView>),
TextureViewArray(Cow<'a, [Arc<TextureView>]>),
AccelerationStructure(Arc<Tlas>),
}
#[derive(Clone, Debug, Error)]

17
wgpu-core/src/command/compute.rs
View File

@ -28,11 +28,11 @@ use crate::{
use thiserror::Error;
use wgt::{BufferAddress, DynamicOffset};
use super::{bind::BinderError, memory_init::CommandBufferTextureMemoryActions};
use crate::ray_tracing::TlasAction;
use std::sync::Arc;
use std::{fmt, mem::size_of, str};
use super::{bind::BinderError, memory_init::CommandBufferTextureMemoryActions};
pub struct ComputePass {
/// All pass data & records is stored here.
///
@ -210,6 +210,7 @@ struct State<'scope, 'snatch_guard, 'cmd_buf, 'raw_encoder> {
tracker: &'cmd_buf mut Tracker,
buffer_memory_init_actions: &'cmd_buf mut Vec<BufferInitTrackerAction>,
texture_memory_actions: &'cmd_buf mut CommandBufferTextureMemoryActions,
tlas_actions: &'cmd_buf mut Vec<TlasAction>,
temp_offsets: Vec<u32>,
dynamic_offset_count: usize,
@ -450,6 +451,7 @@ impl Global {
tracker: &mut cmd_buf_data.trackers,
buffer_memory_init_actions: &mut cmd_buf_data.buffer_memory_init_actions,
texture_memory_actions: &mut cmd_buf_data.texture_memory_actions,
tlas_actions: &mut cmd_buf_data.tlas_actions,
temp_offsets: Vec::new(),
dynamic_offset_count: 0,
@ -695,6 +697,17 @@ fn set_bind_group(
.extend(state.texture_memory_actions.register_init_action(action));
}
let used_resource = bind_group
.used
.acceleration_structures
.into_iter()
.map(|tlas| TlasAction {
tlas: tlas.clone(),
kind: crate::ray_tracing::TlasActionKind::Use,
});
state.tlas_actions.extend(used_resource);
let pipeline_layout = state.binder.pipeline_layout.clone();
let entries = state
.binder

10
wgpu-core/src/command/mod.rs
View File

@ -7,6 +7,7 @@ mod compute_command;
mod draw;
mod memory_init;
mod query;
mod ray_tracing;
mod render;
mod render_command;
mod timestamp_writes;
@ -31,6 +32,7 @@ use crate::lock::{rank, Mutex};
use crate::snatch::SnatchGuard;
use crate::init_tracker::BufferInitTrackerAction;
use crate::ray_tracing::{BlasAction, TlasAction};
use crate::resource::{InvalidResourceError, Labeled};
use crate::track::{DeviceTracker, Tracker, UsageScope};
use crate::LabelHelpers;
@ -242,6 +244,8 @@ impl CommandEncoder {
}
}
/// Look at the documentation for [`CommandBufferMutable`] for an explanation of
/// the fields in this struct. This is the "built" counterpart to that type.
pub(crate) struct BakedCommands {
pub(crate) encoder: Box<dyn hal::DynCommandEncoder>,
pub(crate) list: Vec<Box<dyn hal::DynCommandBuffer>>,
@ -274,6 +278,10 @@ pub struct CommandBufferMutable {
texture_memory_actions: CommandBufferTextureMemoryActions,
pub(crate) pending_query_resets: QueryResetMap,
blas_actions: Vec<BlasAction>,
tlas_actions: Vec<TlasAction>,
#[cfg(feature = "trace")]
pub(crate) commands: Option<Vec<TraceCommand>>,
}
@ -456,6 +464,8 @@ impl CommandBuffer {
buffer_memory_init_actions: Default::default(),
texture_memory_actions: Default::default(),
pending_query_resets: QueryResetMap::new(),
blas_actions: Default::default(),
tlas_actions: Default::default(),
#[cfg(feature = "trace")]
commands: if device.trace.lock().is_some() {
Some(Vec::new())

1285
wgpu-core/src/command/ray_tracing.rs Normal file
View File

File diff suppressed because it is too large Load Diff

8
wgpu-core/src/conv.rs
View File

@ -93,6 +93,14 @@ pub fn map_buffer_usage(usage: wgt::BufferUsages) -> hal::BufferUses {
hal::BufferUses::QUERY_RESOLVE,
usage.contains(wgt::BufferUsages::QUERY_RESOLVE),
);
u.set(
hal::BufferUses::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT,
usage.contains(wgt::BufferUsages::BLAS_INPUT),
);
u.set(
hal::BufferUses::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT,
usage.contains(wgt::BufferUsages::TLAS_INPUT),
);
u
}

21
wgpu-core/src/device/global.rs
View File

@ -732,6 +732,7 @@ impl Global {
buffer_storage: &Storage<Fallible<resource::Buffer>>,
sampler_storage: &Storage<Fallible<resource::Sampler>>,
texture_view_storage: &Storage<Fallible<resource::TextureView>>,
tlas_storage: &Storage<Fallible<resource::Tlas>>,
) -> Result<ResolvedBindGroupEntry<'a>, binding_model::CreateBindGroupError>
{
let resolve_buffer = |bb: &BufferBinding| {
@ -757,6 +758,12 @@ impl Global {
.get()
.map_err(binding_model::CreateBindGroupError::from)
};
let resolve_tlas = |id: &id::TlasId| {
tlas_storage
.get(*id)
.get()
.map_err(binding_model::CreateBindGroupError::from)
};
let resource = match e.resource {
BindingResource::Buffer(ref buffer) => {
ResolvedBindingResource::Buffer(resolve_buffer(buffer)?)
@ -788,6 +795,9 @@ impl Global {
.collect::<Result<Vec<_>, _>>()?;
ResolvedBindingResource::TextureViewArray(Cow::Owned(views))
}
BindingResource::AccelerationStructure(ref tlas) => {
ResolvedBindingResource::AccelerationStructure(resolve_tlas(tlas)?)
}
};
Ok(ResolvedBindGroupEntry {
binding: e.binding,
@ -799,9 +809,18 @@ impl Global {
let buffer_guard = hub.buffers.read();
let texture_view_guard = hub.texture_views.read();
let sampler_guard = hub.samplers.read();
let tlas_guard = hub.tlas_s.read();
desc.entries
.iter()
.map(|e| resolve_entry(e, &buffer_guard, &sampler_guard, &texture_view_guard))
.map(|e| {
resolve_entry(
e,
&buffer_guard,
&sampler_guard,
&texture_view_guard,
&tlas_guard,
)
})
.collect::<Result<Vec<_>, _>>()
};
let entries = match entries {

67
wgpu-core/src/device/life.rs
View File

@ -9,6 +9,7 @@ use crate::{
};
use smallvec::SmallVec;
use crate::resource::{Blas, Tlas};
use std::sync::Arc;
use thiserror::Error;
@ -103,6 +104,44 @@ impl ActiveSubmission {
false
}
pub fn contains_blas(&self, blas: &Blas) -> bool {
for encoder in &self.encoders {
// The ownership location of blas's depends on where the command encoder
// came from. If it is the staging command encoder on the queue, it is
// in the pending buffer list. If it came from a user command encoder,
// it is in the tracker.
if encoder.trackers.blas_s.contains(blas) {
return true;
}
if encoder.pending_blas_s.contains_key(&blas.tracker_index()) {
return true;
}
}
false
}
pub fn contains_tlas(&self, tlas: &Tlas) -> bool {
for encoder in &self.encoders {
// The ownership location of tlas's depends on where the command encoder
// came from. If it is the staging command encoder on the queue, it is
// in the pending buffer list. If it came from a user command encoder,
// it is in the tracker.
if encoder.trackers.tlas_s.contains(tlas) {
return true;
}
if encoder.pending_tlas_s.contains_key(&tlas.tracker_index()) {
return true;
}
}
false
}
}
#[derive(Clone, Debug, Error)]
@ -226,6 +265,34 @@ impl LifetimeTracker {
})
}
/// Returns the submission index of the most recent submission that uses the
/// given blas.
pub fn get_blas_latest_submission_index(&self, blas: &Blas) -> Option<SubmissionIndex> {
// We iterate in reverse order, so that we can bail out early as soon
// as we find a hit.
self.active.iter().rev().find_map(|submission| {
if submission.contains_blas(blas) {
Some(submission.index)
} else {
None
}
})
}
/// Returns the submission index of the most recent submission that uses the
/// given tlas.
pub fn get_tlas_latest_submission_index(&self, tlas: &Tlas) -> Option<SubmissionIndex> {
// We iterate in reverse order, so that we can bail out early as soon
// as we find a hit.
self.active.iter().rev().find_map(|submission| {
if submission.contains_tlas(tlas) {
Some(submission.index)
} else {
None
}
})
}
/// Returns the submission index of the most recent submission that uses the
/// given texture.
pub fn get_texture_latest_submission_index(

5
wgpu-core/src/device/mod.rs
View File

@ -22,6 +22,7 @@ pub(crate) mod bgl;
pub mod global;
mod life;
pub mod queue;
pub mod ray_tracing;
pub mod resource;
#[cfg(any(feature = "trace", feature = "replay"))]
pub mod trace;
@ -544,6 +545,10 @@ pub fn create_validator(
Caps::SUBGROUP_BARRIER,
features.intersects(wgt::Features::SUBGROUP_BARRIER),
);
caps.set(
Caps::RAY_QUERY,
features.intersects(wgt::Features::EXPERIMENTAL_RAY_QUERY),
);
caps.set(
Caps::SUBGROUP_VERTEX_STAGE,
features.contains(wgt::Features::SUBGROUP_VERTEX),

47
wgpu-core/src/device/queue.rs
View File

@ -27,6 +27,8 @@ use crate::{
use smallvec::SmallVec;
use crate::resource::{Blas, DestroyedAccelerationStructure, Tlas};
use crate::scratch::ScratchBuffer;
use std::{
iter,
mem::{self, ManuallyDrop},
@ -143,8 +145,10 @@ impl SubmittedWorkDoneClosure {
#[derive(Debug)]
pub enum TempResource {
StagingBuffer(FlushedStagingBuffer),
ScratchBuffer(ScratchBuffer),
DestroyedBuffer(DestroyedBuffer),
DestroyedTexture(DestroyedTexture),
DestroyedAccelerationStructure(DestroyedAccelerationStructure),
}
/// A series of raw [`CommandBuffer`]s that have been submitted to a
@ -161,6 +165,10 @@ pub(crate) struct EncoderInFlight {
pub(crate) pending_buffers: FastHashMap<TrackerIndex, Arc<Buffer>>,
/// These are the textures that have been tracked by `PendingWrites`.
pub(crate) pending_textures: FastHashMap<TrackerIndex, Arc<Texture>>,
/// These are the BLASes that have been tracked by `PendingWrites`.
pub(crate) pending_blas_s: FastHashMap<TrackerIndex, Arc<Blas>>,
/// These are the TLASes that have been tracked by `PendingWrites`.
pub(crate) pending_tlas_s: FastHashMap<TrackerIndex, Arc<Tlas>>,
}
impl EncoderInFlight {
@ -177,6 +185,8 @@ impl EncoderInFlight {
drop(self.trackers);
drop(self.pending_buffers);
drop(self.pending_textures);
drop(self.pending_blas_s);
drop(self.pending_tlas_s);
}
self.raw
}
@ -216,6 +226,8 @@ pub(crate) struct PendingWrites {
temp_resources: Vec<TempResource>,
dst_buffers: FastHashMap<TrackerIndex, Arc<Buffer>>,
dst_textures: FastHashMap<TrackerIndex, Arc<Texture>>,
dst_blas_s: FastHashMap<TrackerIndex, Arc<Blas>>,
dst_tlas_s: FastHashMap<TrackerIndex, Arc<Tlas>>,
}
impl PendingWrites {
@ -226,6 +238,8 @@ impl PendingWrites {
temp_resources: Vec::new(),
dst_buffers: FastHashMap::default(),
dst_textures: FastHashMap::default(),
dst_blas_s: FastHashMap::default(),
dst_tlas_s: FastHashMap::default(),
}
}
@ -258,6 +272,22 @@ impl PendingWrites {
self.dst_textures.contains_key(&texture.tracker_index())
}
pub fn insert_blas(&mut self, blas: &Arc<Blas>) {
self.dst_blas_s.insert(blas.tracker_index(), blas.clone());
}
pub fn insert_tlas(&mut self, tlas: &Arc<Tlas>) {
self.dst_tlas_s.insert(tlas.tracker_index(), tlas.clone());
}
pub fn contains_blas(&mut self, blas: &Arc<Blas>) -> bool {
self.dst_blas_s.contains_key(&blas.tracker_index())
}
pub fn contains_tlas(&mut self, tlas: &Arc<Tlas>) -> bool {
self.dst_tlas_s.contains_key(&tlas.tracker_index())
}
pub fn consume_temp(&mut self, resource: TempResource) {
self.temp_resources.push(resource);
}
@ -276,6 +306,8 @@ impl PendingWrites {
if self.is_recording {
let pending_buffers = mem::take(&mut self.dst_buffers);
let pending_textures = mem::take(&mut self.dst_textures);
let pending_blas_s = mem::take(&mut self.dst_blas_s);
let pending_tlas_s = mem::take(&mut self.dst_tlas_s);
let cmd_buf = unsafe { self.command_encoder.end_encoding() }
.map_err(|e| device.handle_hal_error(e))?;
@ -291,6 +323,8 @@ impl PendingWrites {
trackers: Tracker::new(),
pending_buffers,
pending_textures,
pending_blas_s,
pending_tlas_s,
};
Ok(Some(encoder))
} else {
@ -358,6 +392,10 @@ pub enum QueueSubmitError {
InvalidResource(#[from] InvalidResourceError),
#[error(transparent)]
CommandEncoder(#[from] CommandEncoderError),
#[error(transparent)]
ValidateBlasActionsError(#[from] crate::ray_tracing::ValidateBlasActionsError),
#[error(transparent)]
ValidateTlasActionsError(#[from] crate::ray_tracing::ValidateTlasActionsError),
}
//TODO: move out common parts of write_xxx.
@ -1117,6 +1155,8 @@ impl Queue {
trackers: baked.trackers,
pending_buffers: FastHashMap::default(),
pending_textures: FastHashMap::default(),
pending_blas_s: FastHashMap::default(),
pending_tlas_s: FastHashMap::default(),
});
}
@ -1462,6 +1502,13 @@ fn validate_command_buffer(
}
}
}
if let Err(e) = cmd_buf_data.validate_blas_actions() {
return Err(e.into());
}
if let Err(e) = cmd_buf_data.validate_tlas_actions(snatch_guard) {
return Err(e.into());
}
}
Ok(())
}

355
wgpu-core/src/device/ray_tracing.rs Normal file
View File

@ -0,0 +1,355 @@
use std::mem::ManuallyDrop;
use std::sync::Arc;
#[cfg(feature = "trace")]
use crate::device::trace;
use crate::lock::{rank, Mutex};
use crate::resource::{Fallible, TrackingData};
use crate::snatch::Snatchable;
use crate::weak_vec::WeakVec;
use crate::{
device::{Device, DeviceError},
global::Global,
id::{self, BlasId, TlasId},
lock::RwLock,
ray_tracing::{get_raw_tlas_instance_size, CreateBlasError, CreateTlasError},
resource, LabelHelpers,
};
use hal::AccelerationStructureTriangleIndices;
use wgt::Features;
impl Device {
fn create_blas(
self: &Arc<Self>,
blas_desc: &resource::BlasDescriptor,
sizes: wgt::BlasGeometrySizeDescriptors,
) -> Result<Arc<resource::Blas>, CreateBlasError> {
let size_info = match &sizes {
wgt::BlasGeometrySizeDescriptors::Triangles { descriptors } => {
let mut entries =
Vec::<hal::AccelerationStructureTriangles<dyn hal::DynBuffer>>::with_capacity(
descriptors.len(),
);
for desc in descriptors {
if desc.index_count.is_some() != desc.index_format.is_some() {
return Err(CreateBlasError::MissingIndexData);
}
let indices =
desc.index_count
.map(|count| AccelerationStructureTriangleIndices::<
dyn hal::DynBuffer,
> {
format: desc.index_format.unwrap(),
buffer: None,
offset: 0,
count,
});
if !self
.features
.allowed_vertex_formats_for_blas()
.contains(&desc.vertex_format)
{
return Err(CreateBlasError::InvalidVertexFormat(
desc.vertex_format,
self.features.allowed_vertex_formats_for_blas(),
));
}
entries.push(hal::AccelerationStructureTriangles::<dyn hal::DynBuffer> {
vertex_buffer: None,
vertex_format: desc.vertex_format,
first_vertex: 0,
vertex_count: desc.vertex_count,
vertex_stride: 0,
indices,
transform: None,
flags: desc.flags,
});
}
unsafe {
self.raw().get_acceleration_structure_build_sizes(
&hal::GetAccelerationStructureBuildSizesDescriptor {
entries: &hal::AccelerationStructureEntries::Triangles(entries),
flags: blas_desc.flags,
},
)
}
}
};
let raw = unsafe {
self.raw()
.create_acceleration_structure(&hal::AccelerationStructureDescriptor {
label: blas_desc.label.as_deref(),
size: size_info.acceleration_structure_size,
format: hal::AccelerationStructureFormat::BottomLevel,
})
}
.map_err(DeviceError::from_hal)?;
let handle = unsafe {
self.raw()
.get_acceleration_structure_device_address(raw.as_ref())
};
Ok(Arc::new(resource::Blas {
raw: Snatchable::new(raw),
device: self.clone(),
size_info,
sizes,
flags: blas_desc.flags,
update_mode: blas_desc.update_mode,
handle,
label: blas_desc.label.to_string(),
built_index: RwLock::new(rank::BLAS_BUILT_INDEX, None),
tracking_data: TrackingData::new(self.tracker_indices.blas_s.clone()),
}))
}
fn create_tlas(
self: &Arc<Self>,
desc: &resource::TlasDescriptor,
) -> Result<Arc<resource::Tlas>, CreateTlasError> {
let size_info = unsafe {
self.raw().get_acceleration_structure_build_sizes(
&hal::GetAccelerationStructureBuildSizesDescriptor {
entries: &hal::AccelerationStructureEntries::Instances(
hal::AccelerationStructureInstances {
buffer: None,
offset: 0,
count: desc.max_instances,
},
),
flags: desc.flags,
},
)
};
let raw = unsafe {
self.raw()
.create_acceleration_structure(&hal::AccelerationStructureDescriptor {
label: desc.label.as_deref(),
size: size_info.acceleration_structure_size,
format: hal::AccelerationStructureFormat::TopLevel,
})
}
.map_err(DeviceError::from_hal)?;
let instance_buffer_size =
get_raw_tlas_instance_size(self.backend()) * desc.max_instances.max(1) as usize;
let instance_buffer = unsafe {
self.raw().create_buffer(&hal::BufferDescriptor {
label: Some("(wgpu-core) instances_buffer"),
size: instance_buffer_size as u64,
usage: hal::BufferUses::COPY_DST
| hal::BufferUses::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT,
memory_flags: hal::MemoryFlags::PREFER_COHERENT,
})
}
.map_err(DeviceError::from_hal)?;
Ok(Arc::new(resource::Tlas {
raw: Snatchable::new(raw),
device: self.clone(),
size_info,
flags: desc.flags,
update_mode: desc.update_mode,
built_index: RwLock::new(rank::TLAS_BUILT_INDEX, None),
dependencies: RwLock::new(rank::TLAS_DEPENDENCIES, Vec::new()),
instance_buffer: ManuallyDrop::new(instance_buffer),
label: desc.label.to_string(),
max_instance_count: desc.max_instances,
tracking_data: TrackingData::new(self.tracker_indices.tlas_s.clone()),
bind_groups: Mutex::new(rank::TLAS_BIND_GROUPS, WeakVec::new()),
}))
}
}
impl Global {
pub fn device_create_blas(
&self,
device_id: id::DeviceId,
desc: &resource::BlasDescriptor,
sizes: wgt::BlasGeometrySizeDescriptors,
id_in: Option<BlasId>,
) -> (BlasId, Option<u64>, Option<CreateBlasError>) {
profiling::scope!("Device::create_blas");
let hub = &self.hub;
let fid = hub.blas_s.prepare(id_in);
let device_guard = hub.devices.read();
let error = 'error: {
let device = device_guard.get(device_id);
match device.check_is_valid() {
Ok(_) => {}
Err(err) => break 'error CreateBlasError::Device(err),
};
if !device
.features
.contains(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)
{
break 'error CreateBlasError::MissingFeature;
}
#[cfg(feature = "trace")]
if let Some(trace) = device.trace.lock().as_mut() {
trace.add(trace::Action::CreateBlas {
id: fid.id(),
desc: desc.clone(),
sizes: sizes.clone(),
});
}
let blas = match device.create_blas(desc, sizes) {
Ok(blas) => blas,
Err(e) => break 'error e,
};
let handle = blas.handle;
let id = fid.assign(Fallible::Valid(blas.clone()));
log::info!("Created blas {:?} with {:?}", id, desc);
return (id, Some(handle), None);
};
let id = fid.assign(Fallible::Invalid(Arc::new(error.to_string())));
(id, None, Some(error))
}
pub fn device_create_tlas(
&self,
device_id: id::DeviceId,
desc: &resource::TlasDescriptor,
id_in: Option<TlasId>,
) -> (TlasId, Option<CreateTlasError>) {
profiling::scope!("Device::create_tlas");
let hub = &self.hub;
let fid = hub.tlas_s.prepare(id_in);
let device_guard = hub.devices.read();
let error = 'error: {
let device = device_guard.get(device_id);
match device.check_is_valid() {
Ok(_) => {}
Err(e) => break 'error CreateTlasError::Device(e),
}
if !device
.features
.contains(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)
{
break 'error CreateTlasError::MissingFeature;
}
#[cfg(feature = "trace")]
if let Some(trace) = device.trace.lock().as_mut() {
trace.add(trace::Action::CreateTlas {
id: fid.id(),
desc: desc.clone(),
});
}
let tlas = match device.create_tlas(desc) {
Ok(tlas) => tlas,
Err(e) => break 'error e,
};
let id = fid.assign(Fallible::Valid(tlas));
log::info!("Created tlas {:?} with {:?}", id, desc);
return (id, None);
};
let id = fid.assign(Fallible::Invalid(Arc::new(error.to_string())));
(id, Some(error))
}
pub fn blas_destroy(&self, blas_id: BlasId) -> Result<(), resource::DestroyError> {
profiling::scope!("Blas::destroy");
log::info!("Blas::destroy {blas_id:?}");
let hub = &self.hub;
let blas = hub.blas_s.get(blas_id).get()?;
let _device = &blas.device;
#[cfg(feature = "trace")]
if let Some(trace) = _device.trace.lock().as_mut() {
trace.add(trace::Action::FreeBlas(blas_id));
}
blas.destroy()
}
pub fn blas_drop(&self, blas_id: BlasId) {
profiling::scope!("Blas::drop");
log::debug!("blas {:?} is dropped", blas_id);
let hub = &self.hub;
let _blas = match hub.blas_s.remove(blas_id).get() {
Ok(blas) => blas,
Err(_) => {
return;
}
};
#[cfg(feature = "trace")]
{
let mut lock = _blas.device.trace.lock();
if let Some(t) = lock.as_mut() {
t.add(trace::Action::DestroyBlas(blas_id));
}
}
}
pub fn tlas_destroy(&self, tlas_id: TlasId) -> Result<(), resource::DestroyError> {
profiling::scope!("Tlas::destroy");
let hub = &self.hub;
log::info!("Tlas {:?} is destroyed", tlas_id);
let tlas_guard = hub.tlas_s.write();
let tlas = tlas_guard
.get(tlas_id)
.get()
.map_err(resource::DestroyError::InvalidResource)?
.clone();
drop(tlas_guard);
let _device = &mut tlas.device.clone();
#[cfg(feature = "trace")]
if let Some(trace) = _device.trace.lock().as_mut() {
trace.add(trace::Action::FreeTlas(tlas_id));
}
tlas.destroy()
}
pub fn tlas_drop(&self, tlas_id: TlasId) {
profiling::scope!("Tlas::drop");
log::debug!("tlas {:?} is dropped", tlas_id);
let hub = &self.hub;
let _tlas = match hub.tlas_s.remove(tlas_id).get() {
Ok(tlas) => tlas,
Err(_) => {
return;
}
};
#[cfg(feature = "trace")]
{
let mut lock = _tlas.device.trace.lock();
if let Some(t) = lock.as_mut() {
t.add(trace::Action::DestroyTlas(tlas_id));
}
}
}
}

46
wgpu-core/src/device/resource.rs
View File

@ -41,6 +41,7 @@ use wgt::{
math::align_to, DeviceLostReason, TextureFormat, TextureSampleType, TextureViewDimension,
};
use crate::resource::{AccelerationStructure, DestroyedResourceError, Tlas};
use std::{
borrow::Cow,
mem::{self, ManuallyDrop},
@ -145,6 +146,7 @@ pub struct Device {
#[cfg(feature = "trace")]
pub(crate) trace: Mutex<Option<trace::Trace>>,
pub(crate) usage_scopes: UsageScopePool,
pub(crate) last_acceleration_structure_build_command_index: AtomicU64,
#[cfg(feature = "indirect-validation")]
pub(crate) indirect_validation: Option<crate::indirect_validation::IndirectValidation>,
@ -334,6 +336,8 @@ impl Device {
),
deferred_destroy: Mutex::new(rank::DEVICE_DEFERRED_DESTROY, Vec::new()),
usage_scopes: Mutex::new(rank::DEVICE_USAGE_SCOPES, Default::default()),
// By starting at one, we can put the result in a NonZeroU64.
last_acceleration_structure_build_command_index: AtomicU64::new(1),
#[cfg(feature = "indirect-validation")]
indirect_validation,
})
@ -1861,7 +1865,7 @@ impl Device {
},
)
}
Bt::AccelerationStructure => todo!(),
Bt::AccelerationStructure => (None, WritableStorage::No),
};
// Validate the count parameter
@ -2188,6 +2192,36 @@ impl Device {
})
}
fn create_tlas_binding<'a>(
self: &Arc<Self>,
used: &mut BindGroupStates,
binding: u32,
decl: &wgt::BindGroupLayoutEntry,
tlas: &'a Arc<Tlas>,
snatch_guard: &'a SnatchGuard<'a>,
) -> Result<&'a dyn hal::DynAccelerationStructure, binding_model::CreateBindGroupError> {
use crate::binding_model::CreateBindGroupError as Error;
used.acceleration_structures.insert_single(tlas.clone());
tlas.same_device(self)?;
match decl.ty {
wgt::BindingType::AccelerationStructure => (),
_ => {
return Err(Error::WrongBindingType {
binding,
actual: decl.ty,
expected: "Tlas",
});
}
}
Ok(tlas
.raw(snatch_guard)
.ok_or(DestroyedResourceError(tlas.error_ident()))?)
}
// This function expects the provided bind group layout to be resolved
// (not passing a duplicate) beforehand.
pub(crate) fn create_bind_group(
@ -2227,6 +2261,7 @@ impl Device {
let mut hal_buffers = Vec::new();
let mut hal_samplers = Vec::new();
let mut hal_textures = Vec::new();
let mut hal_tlas_s = Vec::new();
let snatch_guard = self.snatchable_lock.read();
for entry in desc.entries.iter() {
let binding = entry.binding;
@ -2326,6 +2361,13 @@ impl Device {
(res_index, num_bindings)
}
Br::AccelerationStructure(ref tlas) => {
let tlas =
self.create_tlas_binding(&mut used, binding, decl, tlas, &snatch_guard)?;
let res_index = hal_tlas_s.len();
hal_tlas_s.push(tlas);
(res_index, 1)
}
};
hal_entries.push(hal::BindGroupEntry {
@ -2350,7 +2392,7 @@ impl Device {
buffers: &hal_buffers,
samplers: &hal_samplers,
textures: &hal_textures,
acceleration_structures: &[],
acceleration_structures: &hal_tlas_s,
};
let raw = unsafe { self.raw().create_bind_group(&hal_desc) }
.map_err(|e| self.handle_hal_error(e))?;

21
wgpu-core/src/device/trace.rs
View File

@ -127,6 +127,19 @@ pub enum Action<'a> {
size: wgt::Extent3d,
},
Submit(crate::SubmissionIndex, Vec<Command>),
CreateBlas {
id: id::BlasId,
desc: crate::resource::BlasDescriptor<'a>,
sizes: wgt::BlasGeometrySizeDescriptors,
},
FreeBlas(id::BlasId),
DestroyBlas(id::BlasId),
CreateTlas {
id: id::TlasId,
desc: crate::resource::TlasDescriptor<'a>,
},
FreeTlas(id::TlasId),
DestroyTlas(id::TlasId),
}
#[derive(Debug)]
@ -188,6 +201,14 @@ pub enum Command {
timestamp_writes: Option<crate::command::PassTimestampWrites>,
occlusion_query_set_id: Option<id::QuerySetId>,
},
BuildAccelerationStructuresUnsafeTlas {
blas: Vec<crate::ray_tracing::TraceBlasBuildEntry>,
tlas: Vec<crate::ray_tracing::TlasBuildEntry>,
},
BuildAccelerationStructures {
blas: Vec<crate::ray_tracing::TraceBlasBuildEntry>,
tlas: Vec<crate::ray_tracing::TraceTlasPackage>,
},
}
#[cfg(feature = "trace")]

8
wgpu-core/src/hub.rs
View File

@ -107,7 +107,9 @@ use crate::{
instance::Adapter,
pipeline::{ComputePipeline, PipelineCache, RenderPipeline, ShaderModule},
registry::{Registry, RegistryReport},
resource::{Buffer, Fallible, QuerySet, Sampler, StagingBuffer, Texture, TextureView},
resource::{
Blas, Buffer, Fallible, QuerySet, Sampler, StagingBuffer, Texture, TextureView, Tlas,
},
};
use std::{fmt::Debug, sync::Arc};
@ -178,6 +180,8 @@ pub struct Hub {
pub(crate) textures: Registry<Fallible<Texture>>,
pub(crate) texture_views: Registry<Fallible<TextureView>>,
pub(crate) samplers: Registry<Fallible<Sampler>>,
pub(crate) blas_s: Registry<Fallible<Blas>>,
pub(crate) tlas_s: Registry<Fallible<Tlas>>,
}
impl Hub {
@ -201,6 +205,8 @@ impl Hub {
textures: Registry::new(),
texture_views: Registry::new(),
samplers: Registry::new(),
blas_s: Registry::new(),
tlas_s: Registry::new(),
}
}

3
wgpu-core/src/id.rs
View File

@ -261,6 +261,9 @@ ids! {
pub type RenderBundleEncoderId RenderBundleEncoder;
pub type RenderBundleId RenderBundle;
pub type QuerySetId QuerySet;
pub type BlasId Blas;
pub type TlasId Tlas;
pub type TlasInstanceId TlasInstance;
}
// The CommandBuffer type serves both as encoder and

2
wgpu-core/src/lib.rs
View File

@ -76,6 +76,7 @@ pub mod pipeline;
mod pipeline_cache;
mod pool;
pub mod present;
pub mod ray_tracing;
pub mod registry;
pub mod resource;
mod snatch;
@ -86,6 +87,7 @@ mod weak_vec;
// preserve all run-time checks that `wgpu-core` does.
// See <https://github.com/gfx-rs/wgpu/issues/3103>, after which this can be
// made private again.
mod scratch;
pub mod validation;
pub use hal::{api, MAX_BIND_GROUPS, MAX_COLOR_ATTACHMENTS, MAX_VERTEX_BUFFERS};

4
wgpu-core/src/lock/rank.rs
View File

@ -144,6 +144,10 @@ define_lock_ranks! {
rank TEXTURE_BIND_GROUPS "Texture::bind_groups" followed by { }
rank TEXTURE_INITIALIZATION_STATUS "Texture::initialization_status" followed by { }
rank TEXTURE_VIEWS "Texture::views" followed by { }
rank BLAS_BUILT_INDEX "Blas::built_index" followed by { }
rank TLAS_BUILT_INDEX "Tlas::built_index" followed by { }
rank TLAS_DEPENDENCIES "Tlas::dependencies" followed by { }
rank TLAS_BIND_GROUPS "Tlas::bind_groups" followed by { }
#[cfg(test)]
rank PAWN "pawn" followed by { ROOK, BISHOP }

337
wgpu-core/src/ray_tracing.rs Normal file
View File

@ -0,0 +1,337 @@
// Ray tracing
// Major missing optimizations (no api surface changes needed):
// - use custom tracker to track build state
// - no forced rebuilt (build mode deduction)
// - lazy instance buffer allocation
// - maybe share scratch and instance staging buffer allocation
// - partial instance buffer uploads (api surface already designed with this in mind)
// - ([non performance] extract function in build (rust function extraction with guards is a pain))
use crate::{
command::CommandEncoderError,
device::DeviceError,
id::{BlasId, BufferId, TlasId},
resource::CreateBufferError,
};
use std::sync::Arc;
use std::{num::NonZeroU64, slice};
use crate::resource::{Blas, ResourceErrorIdent, Tlas};
use thiserror::Error;
use wgt::{AccelerationStructureGeometryFlags, BufferAddress, IndexFormat, VertexFormat};
#[derive(Clone, Debug, Error)]
pub enum CreateBlasError {
#[error(transparent)]
Device(#[from] DeviceError),
#[error(transparent)]
CreateBufferError(#[from] CreateBufferError),
#[error(
"Only one of 'index_count' and 'index_format' was provided (either provide both or none)"
)]
MissingIndexData,
#[error("Provided format was not within allowed formats. Provided format: {0:?}. Allowed formats: {1:?}")]
InvalidVertexFormat(VertexFormat, Vec<VertexFormat>),
#[error("Features::RAY_TRACING_ACCELERATION_STRUCTURE is not enabled")]
MissingFeature,
}
#[derive(Clone, Debug, Error)]
pub enum CreateTlasError {
#[error(transparent)]
Device(#[from] DeviceError),
#[error(transparent)]
CreateBufferError(#[from] CreateBufferError),
#[error("Features::RAY_TRACING_ACCELERATION_STRUCTURE is not enabled")]
MissingFeature,
#[error("Unimplemented Tlas error: this error is not yet implemented")]
Unimplemented,
}
/// Error encountered while attempting to do a copy on a command encoder.
#[derive(Clone, Debug, Error)]
pub enum BuildAccelerationStructureError {
#[error(transparent)]
Encoder(#[from] CommandEncoderError),
#[error(transparent)]
Device(#[from] DeviceError),
#[error("BufferId is invalid or destroyed")]
InvalidBufferId,
#[error("Buffer {0:?} is invalid or destroyed")]
InvalidBuffer(ResourceErrorIdent),
#[error("Buffer {0:?} is missing `BLAS_INPUT` usage flag")]
MissingBlasInputUsageFlag(ResourceErrorIdent),
#[error(
"Buffer {0:?} size is insufficient for provided size information (size: {1}, required: {2}"
)]
InsufficientBufferSize(ResourceErrorIdent, u64, u64),
#[error("Buffer {0:?} associated offset doesn't align with the index type")]
UnalignedIndexBufferOffset(ResourceErrorIdent),
#[error("Buffer {0:?} associated offset is unaligned")]
UnalignedTransformBufferOffset(ResourceErrorIdent),
#[error("Buffer {0:?} associated index count not divisible by 3 (count: {1}")]
InvalidIndexCount(ResourceErrorIdent, u32),
#[error("Buffer {0:?} associated data contains None")]
MissingAssociatedData(ResourceErrorIdent),
#[error(
"Blas {0:?} build sizes to may be greater than the descriptor at build time specified"
)]
IncompatibleBlasBuildSizes(ResourceErrorIdent),
#[error("Blas {0:?} flags are different, creation flags: {1:?}, provided: {2:?}")]
IncompatibleBlasFlags(
ResourceErrorIdent,
AccelerationStructureGeometryFlags,
AccelerationStructureGeometryFlags,
),
#[error("Blas {0:?} build vertex count is greater than creation count (needs to be less than or equal to), creation: {1:?}, build: {2:?}")]
IncompatibleBlasVertexCount(ResourceErrorIdent, u32, u32),
#[error("Blas {0:?} vertex formats are different, creation format: {1:?}, provided: {2:?}")]
DifferentBlasVertexFormats(ResourceErrorIdent, VertexFormat, VertexFormat),
#[error("Blas {0:?} index count was provided at creation or building, but not the other")]
BlasIndexCountProvidedMismatch(ResourceErrorIdent),
#[error("Blas {0:?} build index count is greater than creation count (needs to be less than or equal to), creation: {1:?}, build: {2:?}")]
IncompatibleBlasIndexCount(ResourceErrorIdent, u32, u32),
#[error("Blas {0:?} index formats are different, creation format: {1:?}, provided: {2:?}")]
DifferentBlasIndexFormats(ResourceErrorIdent, Option<IndexFormat>, Option<IndexFormat>),
#[error("Blas {0:?} build sizes require index buffer but none was provided")]
MissingIndexBuffer(ResourceErrorIdent),
#[error("BlasId is invalid")]
InvalidBlasId,
#[error("Blas {0:?} is destroyed")]
InvalidBlas(ResourceErrorIdent),
#[error(
"Tlas {0:?} an associated instances contains an invalid custom index (more than 24bits)"
)]
TlasInvalidCustomIndex(ResourceErrorIdent),
#[error(
"Tlas {0:?} has {1} active instances but only {2} are allowed as specified by the descriptor at creation"
)]
TlasInstanceCountExceeded(ResourceErrorIdent, u32, u32),
#[error("BlasId is invalid or destroyed (for instance)")]
InvalidBlasIdForInstance,
#[error("Blas {0:?} is invalid or destroyed (for instance)")]
InvalidBlasForInstance(ResourceErrorIdent),
#[error("TlasId is invalid or destroyed")]
InvalidTlasId,
#[error("Tlas {0:?} is invalid or destroyed")]
InvalidTlas(ResourceErrorIdent),
#[error("Features::RAY_TRACING_ACCELERATION_STRUCTURE is not enabled")]
MissingFeature,
#[error("Buffer {0:?} is missing `TLAS_INPUT` usage flag")]
MissingTlasInputUsageFlag(ResourceErrorIdent),
}
#[derive(Clone, Debug, Error)]
pub enum ValidateBlasActionsError {
#[error("BlasId is invalid or destroyed")]
InvalidBlas,
#[error("Blas {0:?} is used before it is built")]
UsedUnbuilt(ResourceErrorIdent),
}
#[derive(Clone, Debug, Error)]
pub enum ValidateTlasActionsError {
#[error("Tlas {0:?} is invalid or destroyed")]
InvalidTlas(ResourceErrorIdent),
#[error("Tlas {0:?} is used before it is built")]
UsedUnbuilt(ResourceErrorIdent),
#[error("Blas {0:?} is used before it is built (in Tlas {1:?})")]
UsedUnbuiltBlas(ResourceErrorIdent, ResourceErrorIdent),
#[error("BlasId is destroyed (in Tlas {0:?})")]
InvalidBlas(ResourceErrorIdent),
#[error("Blas {0:?} is newer than the containing Tlas {1:?}")]
BlasNewerThenTlas(ResourceErrorIdent, ResourceErrorIdent),
}
#[derive(Debug)]
pub struct BlasTriangleGeometry<'a> {
pub size: &'a wgt::BlasTriangleGeometrySizeDescriptor,
pub vertex_buffer: BufferId,
pub index_buffer: Option<BufferId>,
pub transform_buffer: Option<BufferId>,
pub first_vertex: u32,
pub vertex_stride: BufferAddress,
pub index_buffer_offset: Option<BufferAddress>,
pub transform_buffer_offset: Option<BufferAddress>,
}
pub enum BlasGeometries<'a> {
TriangleGeometries(Box<dyn Iterator<Item = BlasTriangleGeometry<'a>> + 'a>),
}
pub struct BlasBuildEntry<'a> {
pub blas_id: BlasId,
pub geometries: BlasGeometries<'a>,
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct TlasBuildEntry {
pub tlas_id: TlasId,
pub instance_buffer_id: BufferId,
pub instance_count: u32,
}
#[derive(Debug)]
pub struct TlasInstance<'a> {
pub blas_id: BlasId,
pub transform: &'a [f32; 12],
pub custom_index: u32,
pub mask: u8,
}
pub struct TlasPackage<'a> {
pub tlas_id: TlasId,
pub instances: Box<dyn Iterator<Item = Option<TlasInstance<'a>>> + 'a>,
pub lowest_unmodified: u32,
}
#[derive(Debug, Copy, Clone)]
pub(crate) enum BlasActionKind {
Build(NonZeroU64),
Use,
}
#[derive(Debug, Clone)]
pub(crate) enum TlasActionKind {
Build {
build_index: NonZeroU64,
dependencies: Vec<Arc<Blas>>,
},
Use,
}
#[derive(Debug, Clone)]
pub(crate) struct BlasAction {
pub blas: Arc<Blas>,
pub kind: BlasActionKind,
}
#[derive(Debug, Clone)]
pub(crate) struct TlasAction {
pub tlas: Arc<Tlas>,
pub kind: TlasActionKind,
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct TraceBlasTriangleGeometry {
pub size: wgt::BlasTriangleGeometrySizeDescriptor,
pub vertex_buffer: BufferId,
pub index_buffer: Option<BufferId>,
pub transform_buffer: Option<BufferId>,
pub first_vertex: u32,
pub vertex_stride: BufferAddress,
pub index_buffer_offset: Option<BufferAddress>,
pub transform_buffer_offset: Option<BufferAddress>,
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub enum TraceBlasGeometries {
TriangleGeometries(Vec<TraceBlasTriangleGeometry>),
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct TraceBlasBuildEntry {
pub blas_id: BlasId,
pub geometries: TraceBlasGeometries,
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct TraceTlasInstance {
pub blas_id: BlasId,
pub transform: [f32; 12],
pub custom_index: u32,
pub mask: u8,
}
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
pub struct TraceTlasPackage {
pub tlas_id: TlasId,
pub instances: Vec<Option<TraceTlasInstance>>,
pub lowest_unmodified: u32,
}
pub(crate) fn get_raw_tlas_instance_size(backend: wgt::Backend) -> usize {
// TODO: this should be provided by the backend
match backend {
wgt::Backend::Empty => 0,
wgt::Backend::Vulkan => 64,
_ => unimplemented!(),
}
}
#[derive(Clone)]
#[repr(C)]
struct RawTlasInstance {
transform: [f32; 12],
custom_index_and_mask: u32,
shader_binding_table_record_offset_and_flags: u32,
acceleration_structure_reference: u64,
}
pub(crate) fn tlas_instance_into_bytes(
instance: &TlasInstance,
blas_address: u64,
backend: wgt::Backend,
) -> Vec<u8> {
// TODO: get the device to do this
match backend {
wgt::Backend::Empty => vec![],
wgt::Backend::Vulkan => {
const MAX_U24: u32 = (1u32 << 24u32) - 1u32;
let temp = RawTlasInstance {
transform: *instance.transform,
custom_index_and_mask: (instance.custom_index & MAX_U24)
| (u32::from(instance.mask) << 24),
shader_binding_table_record_offset_and_flags: 0,
acceleration_structure_reference: blas_address,
};
let temp: *const _ = &temp;
unsafe {
slice::from_raw_parts::<u8>(
temp.cast::<u8>(),
std::mem::size_of::<RawTlasInstance>(),
)
.to_vec()
}
}
_ => unimplemented!(),
}
}

199
wgpu-core/src/resource.rs
View File

@ -21,6 +21,7 @@ use crate::{
use smallvec::SmallVec;
use thiserror::Error;
use std::num::NonZeroU64;
use std::{
borrow::{Borrow, Cow},
fmt::Debug,
@ -1894,3 +1895,201 @@ pub enum DestroyError {
#[error(transparent)]
InvalidResource(#[from] InvalidResourceError),
}
pub type BlasDescriptor<'a> = wgt::CreateBlasDescriptor<Label<'a>>;
pub type TlasDescriptor<'a> = wgt::CreateTlasDescriptor<Label<'a>>;
pub(crate) trait AccelerationStructure: Trackable {
fn raw<'a>(&'a self, guard: &'a SnatchGuard) -> Option<&'a dyn hal::DynAccelerationStructure>;
}
#[derive(Debug)]
pub struct Blas {
pub(crate) raw: Snatchable<Box<dyn hal::DynAccelerationStructure>>,
pub(crate) device: Arc<Device>,
pub(crate) size_info: hal::AccelerationStructureBuildSizes,
pub(crate) sizes: wgt::BlasGeometrySizeDescriptors,
pub(crate) flags: wgt::AccelerationStructureFlags,
pub(crate) update_mode: wgt::AccelerationStructureUpdateMode,
pub(crate) built_index: RwLock<Option<NonZeroU64>>,
pub(crate) handle: u64,
/// The `label` from the descriptor used to create the resource.
pub(crate) label: String,
pub(crate) tracking_data: TrackingData,
}
impl Drop for Blas {
fn drop(&mut self) {
resource_log!("Destroy raw {}", self.error_ident());
// SAFETY: We are in the Drop impl, and we don't use self.raw anymore after this point.
if let Some(raw) = self.raw.take() {
unsafe {
self.device.raw().destroy_acceleration_structure(raw);
}
}
}
}
impl AccelerationStructure for Blas {
fn raw<'a>(&'a self, guard: &'a SnatchGuard) -> Option<&'a dyn hal::DynAccelerationStructure> {
Some(self.raw.get(guard)?.as_ref())
}
}
impl Blas {
pub(crate) fn destroy(self: &Arc<Self>) -> Result<(), DestroyError> {
let device = &self.device;
let temp = {
let mut snatch_guard = device.snatchable_lock.write();
let raw = match self.raw.snatch(&mut snatch_guard) {
Some(raw) => raw,
None => {
return Err(DestroyError::AlreadyDestroyed);
}
};
drop(snatch_guard);
queue::TempResource::DestroyedAccelerationStructure(DestroyedAccelerationStructure {
raw: ManuallyDrop::new(raw),
device: Arc::clone(&self.device),
label: self.label().to_owned(),
bind_groups: WeakVec::new(),
})
};
let mut pending_writes = device.pending_writes.lock();
if pending_writes.contains_blas(self) {
pending_writes.consume_temp(temp);
} else {
let mut life_lock = device.lock_life();
let last_submit_index = life_lock.get_blas_latest_submission_index(self);
if let Some(last_submit_index) = last_submit_index {
life_lock.schedule_resource_destruction(temp, last_submit_index);
}
}
Ok(())
}
}
crate::impl_resource_type!(Blas);
crate::impl_labeled!(Blas);
crate::impl_parent_device!(Blas);
crate::impl_storage_item!(Blas);
crate::impl_trackable!(Blas);
#[derive(Debug)]
pub struct Tlas {
pub(crate) raw: Snatchable<Box<dyn hal::DynAccelerationStructure>>,
pub(crate) device: Arc<Device>,
pub(crate) size_info: hal::AccelerationStructureBuildSizes,
pub(crate) max_instance_count: u32,
pub(crate) flags: wgt::AccelerationStructureFlags,
pub(crate) update_mode: wgt::AccelerationStructureUpdateMode,
pub(crate) built_index: RwLock<Option<NonZeroU64>>,
pub(crate) dependencies: RwLock<Vec<Arc<Blas>>>,
pub(crate) instance_buffer: ManuallyDrop<Box<dyn hal::DynBuffer>>,
/// The `label` from the descriptor used to create the resource.
pub(crate) label: String,
pub(crate) tracking_data: TrackingData,
pub(crate) bind_groups: Mutex<WeakVec<BindGroup>>,
}
impl Drop for Tlas {
fn drop(&mut self) {
unsafe {
resource_log!("Destroy raw {}", self.error_ident());
if let Some(structure) = self.raw.take() {
self.device.raw().destroy_acceleration_structure(structure);
}
let buffer = ManuallyDrop::take(&mut self.instance_buffer);
self.device.raw().destroy_buffer(buffer);
}
}
}
impl AccelerationStructure for Tlas {
fn raw<'a>(&'a self, guard: &'a SnatchGuard) -> Option<&dyn hal::DynAccelerationStructure> {
Some(self.raw.get(guard)?.as_ref())
}
}
crate::impl_resource_type!(Tlas);
crate::impl_labeled!(Tlas);
crate::impl_parent_device!(Tlas);
crate::impl_storage_item!(Tlas);
crate::impl_trackable!(Tlas);
impl Tlas {
pub(crate) fn destroy(self: &Arc<Self>) -> Result<(), DestroyError> {
let device = &self.device;
let temp = {
let mut snatch_guard = device.snatchable_lock.write();
let raw = match self.raw.snatch(&mut snatch_guard) {
Some(raw) => raw,
None => {
return Err(DestroyError::AlreadyDestroyed);
}
};
drop(snatch_guard);
queue::TempResource::DestroyedAccelerationStructure(DestroyedAccelerationStructure {
raw: ManuallyDrop::new(raw),
device: Arc::clone(&self.device),
label: self.label().to_owned(),
bind_groups: mem::take(&mut self.bind_groups.lock()),
})
};
let mut pending_writes = device.pending_writes.lock();
if pending_writes.contains_tlas(self) {
pending_writes.consume_temp(temp);
} else {
let mut life_lock = device.lock_life();
let last_submit_index = life_lock.get_tlas_latest_submission_index(self);
if let Some(last_submit_index) = last_submit_index {
life_lock.schedule_resource_destruction(temp, last_submit_index);
}
}
Ok(())
}
}
#[derive(Debug)]
pub struct DestroyedAccelerationStructure {
raw: ManuallyDrop<Box<dyn hal::DynAccelerationStructure>>,
device: Arc<Device>,
label: String,
// only filled if the acceleration structure is a TLAS
bind_groups: WeakVec<BindGroup>,
}
impl DestroyedAccelerationStructure {
pub fn label(&self) -> &dyn Debug {
&self.label
}
}
impl Drop for DestroyedAccelerationStructure {
fn drop(&mut self) {
let mut deferred = self.device.deferred_destroy.lock();
deferred.push(DeferredDestroy::BindGroups(mem::take(
&mut self.bind_groups,
)));
drop(deferred);
resource_log!("Destroy raw Buffer (destroyed) {:?}", self.label());
// SAFETY: We are in the Drop impl and we don't use self.raw anymore after this point.
let raw = unsafe { ManuallyDrop::take(&mut self.raw) };
unsafe {
hal::DynDevice::destroy_acceleration_structure(self.device.raw(), raw);
}
}
}

43
wgpu-core/src/scratch.rs Normal file
View File

@ -0,0 +1,43 @@
use crate::device::{Device, DeviceError};
use crate::resource_log;
use hal::BufferUses;
use std::mem::ManuallyDrop;
use std::sync::Arc;
#[derive(Debug)]
pub struct ScratchBuffer {
raw: ManuallyDrop<Box<dyn hal::DynBuffer>>,
device: Arc<Device>,
}
impl ScratchBuffer {
pub(crate) fn new(device: &Arc<Device>, size: wgt::BufferSize) -> Result<Self, DeviceError> {
let raw = unsafe {
device
.raw()
.create_buffer(&hal::BufferDescriptor {
label: Some("(wgpu) scratch buffer"),
size: size.get(),
usage: BufferUses::ACCELERATION_STRUCTURE_SCRATCH,
memory_flags: hal::MemoryFlags::empty(),
})
.map_err(crate::device::DeviceError::from_hal)?
};
Ok(Self {
raw: ManuallyDrop::new(raw),
device: device.clone(),
})
}
pub(crate) fn raw(&self) -> &dyn hal::DynBuffer {
self.raw.as_ref()
}
}
impl Drop for ScratchBuffer {
fn drop(&mut self) {
resource_log!("Destroy raw ScratchBuffer");
// SAFETY: We are in the Drop impl and we don't use self.raw anymore after this point.
let raw = unsafe { ManuallyDrop::take(&mut self.raw) };
unsafe { self.device.raw().destroy_buffer(raw) };
}
}

14
wgpu-core/src/track/mod.rs
View File

@ -98,6 +98,7 @@ Device <- CommandBuffer = insert(device.start, device.end, buffer.start, buffer.
mod buffer;
mod metadata;
mod range;
mod ray_tracing;
mod stateless;
mod texture;
@ -112,6 +113,7 @@ use crate::{
use std::{fmt, ops, sync::Arc};
use thiserror::Error;
use crate::track::ray_tracing::AccelerationStructureTracker;
pub(crate) use buffer::{
BufferBindGroupState, BufferTracker, BufferUsageScope, DeviceBufferTracker,
};
@ -224,6 +226,8 @@ pub(crate) struct TrackerIndexAllocators {
pub render_pipelines: Arc<SharedTrackerIndexAllocator>,
pub bundles: Arc<SharedTrackerIndexAllocator>,
pub query_sets: Arc<SharedTrackerIndexAllocator>,
pub blas_s: Arc<SharedTrackerIndexAllocator>,
pub tlas_s: Arc<SharedTrackerIndexAllocator>,
}
impl TrackerIndexAllocators {
@ -238,6 +242,8 @@ impl TrackerIndexAllocators {
render_pipelines: Arc::new(SharedTrackerIndexAllocator::new()),
bundles: Arc::new(SharedTrackerIndexAllocator::new()),
query_sets: Arc::new(SharedTrackerIndexAllocator::new()),
blas_s: Arc::new(SharedTrackerIndexAllocator::new()),
tlas_s: Arc::new(SharedTrackerIndexAllocator::new()),
}
}
}
@ -420,6 +426,7 @@ pub(crate) struct BindGroupStates {
pub buffers: BufferBindGroupState,
pub views: TextureViewBindGroupState,
pub samplers: StatelessTracker<resource::Sampler>,
pub acceleration_structures: StatelessTracker<resource::Tlas>,
}
impl BindGroupStates {
@ -428,6 +435,7 @@ impl BindGroupStates {
buffers: BufferBindGroupState::new(),
views: TextureViewBindGroupState::new(),
samplers: StatelessTracker::new(),
acceleration_structures: StatelessTracker::new(),
}
}
@ -440,7 +448,7 @@ impl BindGroupStates {
// Views are stateless, however, `TextureViewBindGroupState`
// is special as it will be merged with other texture trackers.
self.views.optimize();
// Samplers are stateless and don't need to be optimized
// Samplers and Tlas's are stateless and don't need to be optimized
// since the tracker is never merged with any other tracker.
}
}
@ -594,6 +602,8 @@ impl DeviceTracker {
pub(crate) struct Tracker {
pub buffers: BufferTracker,
pub textures: TextureTracker,
pub blas_s: AccelerationStructureTracker<resource::Blas>,
pub tlas_s: AccelerationStructureTracker<resource::Tlas>,
pub views: StatelessTracker<resource::TextureView>,
pub bind_groups: StatelessTracker<binding_model::BindGroup>,
pub compute_pipelines: StatelessTracker<pipeline::ComputePipeline>,
@ -607,6 +617,8 @@ impl Tracker {
Self {
buffers: BufferTracker::new(),
textures: TextureTracker::new(),
blas_s: AccelerationStructureTracker::new(),
tlas_s: AccelerationStructureTracker::new(),
views: StatelessTracker::new(),
bind_groups: StatelessTracker::new(),
compute_pipelines: StatelessTracker::new(),

81
wgpu-core/src/track/ray_tracing.rs Normal file
View File

@ -0,0 +1,81 @@
use crate::resource::AccelerationStructure;
use crate::track::metadata::ResourceMetadata;
use crate::track::ResourceUses;
use hal::AccelerationStructureUses;
use std::sync::Arc;
use wgt::strict_assert;
pub(crate) struct AccelerationStructureTracker<T: AccelerationStructure> {
start: Vec<AccelerationStructureUses>,
end: Vec<AccelerationStructureUses>,
metadata: ResourceMetadata<Arc<T>>,
}
impl<T: AccelerationStructure> AccelerationStructureTracker<T> {
pub fn new() -> Self {
Self {
start: Vec::new(),
end: Vec::new(),
metadata: ResourceMetadata::new(),
}
}
fn tracker_assert_in_bounds(&self, index: usize) {
strict_assert!(index < self.start.len());
strict_assert!(index < self.end.len());
self.metadata.tracker_assert_in_bounds(index);
}
/// Sets the size of all the vectors inside the tracker.
///
/// Must be called with the highest possible Buffer ID before
/// all unsafe functions are called.
pub fn set_size(&mut self, size: usize) {
self.start.resize(size, AccelerationStructureUses::empty());
self.end.resize(size, AccelerationStructureUses::empty());
self.metadata.set_size(size);
}
/// Extend the vectors to let the given index be valid.
fn allow_index(&mut self, index: usize) {
if index >= self.start.len() {
self.set_size(index + 1);
}
}
/// Returns true if the given buffer is tracked.
pub fn contains(&self, acceleration_structure: &T) -> bool {
self.metadata
.contains(acceleration_structure.tracker_index().as_usize())
}
/// Inserts a single resource into the resource tracker.
pub fn set_single(&mut self, resource: Arc<T>) {
let index: usize = resource.tracker_index().as_usize();
self.allow_index(index);
self.tracker_assert_in_bounds(index);
}
}
impl ResourceUses for AccelerationStructureUses {
const EXCLUSIVE: Self = Self::empty();
type Selector = ();
fn bits(self) -> u16 {
Self::bits(&self) as u16
}
fn all_ordered(self) -> bool {
true
}
fn any_exclusive(self) -> bool {
self.intersects(Self::EXCLUSIVE)
}
}

10
wgpu-core/src/track/stateless.rs
View File

@ -1,3 +1,4 @@
use std::slice::Iter;
use std::sync::Arc;
/// A tracker that holds strong references to resources.
@ -24,3 +25,12 @@ impl<T> StatelessTracker<T> {
unsafe { self.resources.last().unwrap_unchecked() }
}
}
impl<'a, T> IntoIterator for &'a StatelessTracker<T> {
type Item = &'a Arc<T>;
type IntoIter = Iter<'a, Arc<T>>;
fn into_iter(self) -> Self::IntoIter {
self.resources.as_slice().iter()
}
}

4
wgpu-core/src/validation.rs
View File

@ -17,6 +17,7 @@ enum ResourceType {
Sampler {
comparison: bool,
},
AccelerationStructure,
}
#[derive(Debug)]
@ -490,6 +491,7 @@ impl Resource {
});
}
}
ResourceType::AccelerationStructure => (),
};
Ok(())
@ -567,6 +569,7 @@ impl Resource {
},
}
}
ResourceType::AccelerationStructure => BindingType::AccelerationStructure,
})
}
}
@ -861,6 +864,7 @@ impl Interface {
class,
},
naga::TypeInner::Sampler { comparison } => ResourceType::Sampler { comparison },
naga::TypeInner::AccelerationStructure => ResourceType::AccelerationStructure,
ref other => ResourceType::Buffer {
size: wgt::BufferSize::new(other.size(module.to_ctx()) as u64).unwrap(),
},

24
wgpu-hal/src/vulkan/adapter.rs
View File

@ -77,17 +77,17 @@ pub struct PhysicalDeviceFeatures {
/// Features provided by `VK_KHR_buffer_device_address`, promoted to Vulkan 1.2.
///
/// We only use this feature for
/// [`Features::RAY_TRACING_ACCELERATION_STRUCTURE`], which requires
/// [`Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE`], which requires
/// `VK_KHR_acceleration_structure`, which depends on
/// `VK_KHR_buffer_device_address`, so [`Instance::expose_adapter`] only
/// bothers to check if `VK_KHR_acceleration_structure` is available,
/// leaving this `None`.
///
/// However, we do populate this when creating a device if
/// [`Features::RAY_TRACING_ACCELERATION_STRUCTURE`] is requested.
/// [`Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE`] is requested.
///
/// [`Instance::expose_adapter`]: super::Instance::expose_adapter
/// [`Features::RAY_TRACING_ACCELERATION_STRUCTURE`]: wgt::Features::RAY_TRACING_ACCELERATION_STRUCTURE
/// [`Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE`]: wgt::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
buffer_device_address: Option<vk::PhysicalDeviceBufferDeviceAddressFeaturesKHR<'static>>,
/// Features provided by `VK_KHR_ray_query`,
@ -722,13 +722,16 @@ impl PhysicalDeviceFeatures {
features.set(F::DEPTH32FLOAT_STENCIL8, texture_d32_s8);
features.set(
F::RAY_TRACING_ACCELERATION_STRUCTURE,
F::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE,
caps.supports_extension(khr::deferred_host_operations::NAME)
&& caps.supports_extension(khr::acceleration_structure::NAME)
&& caps.supports_extension(khr::buffer_device_address::NAME),
);
features.set(F::RAY_QUERY, caps.supports_extension(khr::ray_query::NAME));
features.set(
F::EXPERIMENTAL_RAY_QUERY,
caps.supports_extension(khr::ray_query::NAME),
);
let rg11b10ufloat_renderable = supports_format(
instance,
@ -983,14 +986,16 @@ impl PhysicalDeviceProperties {
}
// Require `VK_KHR_deferred_host_operations`, `VK_KHR_acceleration_structure` and `VK_KHR_buffer_device_address` if the feature `RAY_TRACING` was requested
if requested_features.contains(wgt::Features::RAY_TRACING_ACCELERATION_STRUCTURE) {
if requested_features
.contains(wgt::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)
{
extensions.push(khr::deferred_host_operations::NAME);
extensions.push(khr::acceleration_structure::NAME);
extensions.push(khr::buffer_device_address::NAME);
}
// Require `VK_KHR_ray_query` if the associated feature was requested
if requested_features.contains(wgt::Features::RAY_QUERY) {
if requested_features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) {
extensions.push(khr::ray_query::NAME);
}
@ -1779,7 +1784,7 @@ impl super::Adapter {
capabilities.push(spv::Capability::StorageImageWriteWithoutFormat);
}
if features.contains(wgt::Features::RAY_QUERY) {
if features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) {
capabilities.push(spv::Capability::RayQueryKHR);
}
@ -1810,6 +1815,9 @@ impl super::Adapter {
// But this requires cloning the `spv::Options` struct, which has heap allocations.
true, // could check `super::Workarounds::SEPARATE_ENTRY_POINTS`
);
if features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) {
capabilities.push(spv::Capability::RayQueryKHR);
}
spv::Options {
lang_version: if features
.intersects(wgt::Features::SUBGROUP | wgt::Features::SUBGROUP_VERTEX)

12
wgpu-hal/src/vulkan/command.rs
View File

@ -647,10 +647,14 @@ impl crate::CommandEncoder for super::CommandEncoder {
&mut self,
barrier: crate::AccelerationStructureBarrier,
) {
let (src_stage, src_access) =
conv::map_acceleration_structure_usage_to_barrier(barrier.usage.start);
let (dst_stage, dst_access) =
conv::map_acceleration_structure_usage_to_barrier(barrier.usage.end);
let (src_stage, src_access) = conv::map_acceleration_structure_usage_to_barrier(
barrier.usage.start,
self.device.features,
);
let (dst_stage, dst_access) = conv::map_acceleration_structure_usage_to_barrier(
barrier.usage.end,
self.device.features,
);
unsafe {
self.device.raw.cmd_pipeline_barrier(

5
wgpu-hal/src/vulkan/conv.rs
View File

@ -943,6 +943,7 @@ pub fn map_acceleration_structure_geometry_flags(
pub fn map_acceleration_structure_usage_to_barrier(
usage: crate::AccelerationStructureUses,
features: wgt::Features,
) -> (vk::PipelineStageFlags, vk::AccessFlags) {
let mut stages = vk::PipelineStageFlags::empty();
let mut access = vk::AccessFlags::empty();
@ -955,7 +956,9 @@ pub fn map_acceleration_structure_usage_to_barrier(
stages |= vk::PipelineStageFlags::ACCELERATION_STRUCTURE_BUILD_KHR;
access |= vk::AccessFlags::ACCELERATION_STRUCTURE_WRITE_KHR;
}
if usage.contains(crate::AccelerationStructureUses::SHADER_INPUT) {
if usage.contains(crate::AccelerationStructureUses::SHADER_INPUT)
&& features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY)
{
stages |= vk::PipelineStageFlags::VERTEX_SHADER
| vk::PipelineStageFlags::FRAGMENT_SHADER
| vk::PipelineStageFlags::COMPUTE_SHADER;

9
wgpu-hal/src/vulkan/device.rs
View File

@ -1075,14 +1075,7 @@ impl crate::Device for super::Device {
desc.memory_flags.contains(crate::MemoryFlags::TRANSIENT),
);
let alignment_mask = if desc.usage.intersects(
crate::BufferUses::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT
| crate::BufferUses::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT,
) {
16
} else {
req.alignment
} - 1;
let alignment_mask = req.alignment - 1;
let block = unsafe {
self.mem_allocator.lock().alloc(

156
wgpu-types/src/lib.rs
View File

@ -828,23 +828,32 @@ bitflags::bitflags! {
///
/// This is a native only feature.
const TEXTURE_FORMAT_NV12 = 1 << 47;
/// Allows for the creation of ray-tracing acceleration structures.
/// ***THIS IS EXPERIMENTAL:*** Features enabled by this may have
/// major bugs in them and are expected to be subject to breaking changes, suggestions
/// for the API exposed by this should be posted on [the ray-tracing issue](https://github.com/gfx-rs/wgpu/issues/1040)
///
/// Allows for the creation of ray-tracing acceleration structures. Currently,
/// ray-tracing acceleration structures are only useful when used with [Features::EXPERIMENTAL_RAY_QUERY]
///
/// Supported platforms:
/// - Vulkan
///
/// This is a native-only feature.
const RAY_TRACING_ACCELERATION_STRUCTURE = 1 << 48;
const EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE = 1 << 48;
// Shader:
/// ***THIS IS EXPERIMENTAL:*** Features enabled by this may have
/// major bugs in it and are expected to be subject to breaking changes, suggestions
/// for the API exposed by this should be posted on [the ray-tracing issue](https://github.com/gfx-rs/wgpu/issues/1040)
///
/// Allows for the creation of ray-tracing queries within shaders.
///
/// Supported platforms:
/// - Vulkan
///
/// This is a native-only feature.
const RAY_QUERY = 1 << 49;
const EXPERIMENTAL_RAY_QUERY = 1 << 49;
/// Enables 64-bit floating point types in SPIR-V shaders.
///
/// Note: even when supported by GPU hardware, 64-bit floating point operations are
@ -986,6 +995,16 @@ impl Features {
pub const fn all_native_mask() -> Self {
Self::from_bits_truncate(!Self::all_webgpu_mask().bits())
}
/// Vertex formats allowed for creating and building BLASes
#[must_use]
pub fn allowed_vertex_formats_for_blas(&self) -> Vec<VertexFormat> {
let mut formats = Vec::new();
if self.contains(Self::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE) {
formats.push(VertexFormat::Float32x3);
}
formats
}
}
bitflags::bitflags! {
@ -5296,6 +5315,10 @@ bitflags::bitflags! {
const INDIRECT = 1 << 8;
/// Allow a buffer to be the destination buffer for a [`CommandEncoder::resolve_query_set`] operation.
const QUERY_RESOLVE = 1 << 9;
/// Allows a buffer to be used as input for a bottom level acceleration structure build
const BLAS_INPUT = 1 << 10;
/// Allows a buffer to be used as input for a top level acceleration structure build
const TLAS_INPUT = 1 << 11;
}
}
@ -7423,19 +7446,119 @@ impl Default for InstanceDescriptor {
}
}
#[derive(Clone, Debug, PartialEq, Eq)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
/// Descriptor for all size defining attributes of a single triangle geometry inside a bottom level acceleration structure.
pub struct BlasTriangleGeometrySizeDescriptor {
/// Format of a vertex position, must be [VertexFormat::Float32x3]
/// with just [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE]
/// but later features may add more formats.
pub vertex_format: VertexFormat,
/// Number of vertices.
pub vertex_count: u32,
/// Format of an index. Only needed if an index buffer is used.
/// If `index_format` is provided `index_count` is required.
pub index_format: Option<IndexFormat>,
/// Number of indices. Only needed if an index buffer is used.
/// If `index_count` is provided `index_format` is required.
pub index_count: Option<u32>,
/// Flags for the geometry.
pub flags: AccelerationStructureGeometryFlags,
}
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
/// Descriptor for all size defining attributes of all geometries inside a bottom level acceleration structure.
pub enum BlasGeometrySizeDescriptors {
/// Triangle geometry version.
Triangles {
/// Descriptor for each triangle geometry.
descriptors: Vec<BlasTriangleGeometrySizeDescriptor>,
},
}
#[repr(u8)]
#[derive(Copy, Clone, Debug, PartialEq, Eq, Hash)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
/// Update mode for acceleration structure builds.
pub enum AccelerationStructureUpdateMode {
/// Always perform a full build.
Build,
/// If possible, perform an incremental update.
///
/// Not advised for major topology changes.
/// (Useful for e.g. skinning)
PreferUpdate,
}
#[repr(C)]
#[derive(Clone, Debug, PartialEq, Eq, Hash)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
/// Descriptor for creating a bottom level acceleration structure.
pub struct CreateBlasDescriptor<L> {
/// Label for the bottom level acceleration structure.
pub label: L,
/// Flags for the bottom level acceleration structure.
pub flags: AccelerationStructureFlags,
/// Update mode for the bottom level acceleration structure.
pub update_mode: AccelerationStructureUpdateMode,
}
impl<L> CreateBlasDescriptor<L> {
/// Takes a closure and maps the label of the blas descriptor into another.
pub fn map_label<K>(&self, fun: impl FnOnce(&L) -> K) -> CreateBlasDescriptor<K> {
CreateBlasDescriptor {
label: fun(&self.label),
flags: self.flags,
update_mode: self.update_mode,
}
}
}
#[repr(C)]
#[derive(Clone, Debug, PartialEq, Eq, Hash)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
/// Descriptor for creating a top level acceleration structure.
pub struct CreateTlasDescriptor<L> {
/// Label for the top level acceleration structure.
pub label: L,
/// Number of instances that can be stored in the acceleration structure.
pub max_instances: u32,
/// Flags for the bottom level acceleration structure.
pub flags: AccelerationStructureFlags,
/// Update mode for the bottom level acceleration structure.
pub update_mode: AccelerationStructureUpdateMode,
}
impl<L> CreateTlasDescriptor<L> {
/// Takes a closure and maps the label of the blas descriptor into another.
pub fn map_label<K>(&self, fun: impl FnOnce(&L) -> K) -> CreateTlasDescriptor<K> {
CreateTlasDescriptor {
label: fun(&self.label),
flags: self.flags,
update_mode: self.update_mode,
max_instances: self.max_instances,
}
}
}
bitflags::bitflags!(
/// Flags for acceleration structures
#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
pub struct AccelerationStructureFlags: u8 {
/// Allow for incremental updates (no change in size)
/// Allow for incremental updates (no change in size), currently this is unimplemented
/// and will build as normal (this is fine, update vs build should be unnoticeable)
const ALLOW_UPDATE = 1 << 0;
/// Allow the acceleration structure to be compacted in a copy operation
/// Allow the acceleration structure to be compacted in a copy operation, the function
/// to compact is not currently implemented.
const ALLOW_COMPACTION = 1 << 1;
/// Optimize for fast ray tracing performance
/// Optimize for fast ray tracing performance, recommended if the geometry is unlikely
/// to change (e.g. in a game: non-interactive scene geometry)
const PREFER_FAST_TRACE = 1 << 2;
/// Optimize for fast build time
/// Optimize for fast build time, recommended if geometry is likely to change frequently
/// (e.g. in a game: player model).
const PREFER_FAST_BUILD = 1 << 3;
/// Optimize for low memory footprint (scratch and output)
/// Optimize for low memory footprint (both while building and in the output BLAS).
const LOW_MEMORY = 1 << 4;
}
);
@ -7445,14 +7568,27 @@ bitflags::bitflags!(
/// Flags for acceleration structure geometries
#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
pub struct AccelerationStructureGeometryFlags: u8 {
/// Is OPAQUE
/// Is OPAQUE (is there no alpha test) recommended as currently in naga there is no
/// candidate intersections yet so currently BLASes without this flag will not have hits.
/// Not enabling this makes the BLAS unable to be interacted with in WGSL.
const OPAQUE = 1 << 0;
/// NO_DUPLICATE_ANY_HIT_INVOCATION
/// NO_DUPLICATE_ANY_HIT_INVOCATION, not useful unless using hal with wgpu, ray-tracing
/// pipelines are not supported in wgpu so any-hit shaders do not exist. For when any-hit
/// shaders are implemented (or experienced users who combine this with an underlying library:
/// for any primitive (triangle or AABB) multiple any-hit shaders sometimes may be invoked
/// (especially in AABBs like a sphere), if this flag in present only one hit on a primitive may
/// invoke an any-hit shader.
const NO_DUPLICATE_ANY_HIT_INVOCATION = 1 << 1;
}
);
impl_bitflags!(AccelerationStructureGeometryFlags);
/// Alignment requirement for transform buffers used in acceleration structure builds
pub const TRANSFORM_BUFFER_ALIGNMENT: BufferAddress = 16;
/// Alignment requirement for instance buffers used in acceleration structure builds (`build_acceleration_structures_unsafe_tlas`)
pub const INSTANCE_BUFFER_ALIGNMENT: BufferAddress = 16;
pub use send_sync::*;
#[doc(hidden)]

12
wgpu/src/api/bind_group.rs
View File

@ -70,6 +70,18 @@ pub enum BindingResource<'a> {
/// Corresponds to [`wgt::BindingType::Texture`] and [`wgt::BindingType::StorageTexture`] with
/// [`BindGroupLayoutEntry::count`] set to Some.
TextureViewArray(&'a [&'a TextureView]),
/// Binding is backed by a top level acceleration structure
///
/// Corresponds to [`wgt::BindingType::AccelerationStructure`] with [`BindGroupLayoutEntry::count`] set to None.
///
/// # Validation
/// When using (e.g. with `set_bind_group`) a bind group that has been created with one or more of this binding
/// resource certain checks take place.
/// - TLAS must have been built, if not a validation error is generated
/// - All BLASes that were built into the TLAS must be built before the TLAS, if this was not satisfied and TLAS was
/// built using `build_acceleration_structures` a validation error is generated otherwise this is a part of the
/// safety section of `build_acceleration_structures_unsafe_tlas` and so undefined behavior occurs.
AccelerationStructure(&'a Tlas),
}
#[cfg(send_sync)]
static_assertions::assert_impl_all!(BindingResource<'_>: Send, Sync);

232
wgpu/src/api/blas.rs Normal file
View File

@ -0,0 +1,232 @@
use crate::context::{Context, DynContext};
use crate::{Buffer, Data, Label, C};
use std::sync::Arc;
use std::thread;
use wgt::WasmNotSendSync;
/// Descriptor for the size defining attributes of a triangle geometry, for a bottom level acceleration structure.
pub type BlasTriangleGeometrySizeDescriptor = wgt::BlasTriangleGeometrySizeDescriptor;
static_assertions::assert_impl_all!(BlasTriangleGeometrySizeDescriptor: Send, Sync);
/// Descriptor for the size defining attributes, for a bottom level acceleration structure.
pub type BlasGeometrySizeDescriptors = wgt::BlasGeometrySizeDescriptors;
static_assertions::assert_impl_all!(BlasGeometrySizeDescriptors: Send, Sync);
/// Flags for an acceleration structure.
pub type AccelerationStructureFlags = wgt::AccelerationStructureFlags;
static_assertions::assert_impl_all!(AccelerationStructureFlags: Send, Sync);
/// Flags for a geometry inside a bottom level acceleration structure.
pub type AccelerationStructureGeometryFlags = wgt::AccelerationStructureGeometryFlags;
static_assertions::assert_impl_all!(AccelerationStructureGeometryFlags: Send, Sync);
/// Update mode for acceleration structure builds.
pub type AccelerationStructureUpdateMode = wgt::AccelerationStructureUpdateMode;
static_assertions::assert_impl_all!(AccelerationStructureUpdateMode: Send, Sync);
/// Descriptor to create bottom level acceleration structures.
pub type CreateBlasDescriptor<'a> = wgt::CreateBlasDescriptor<Label<'a>>;
static_assertions::assert_impl_all!(CreateBlasDescriptor<'_>: Send, Sync);
/// Safe instance for a [Tlas].
///
/// A TlasInstance may be made invalid, if a TlasInstance is invalid, any attempt to build a [TlasPackage] containing an
/// invalid TlasInstance will generate a validation error
///
/// Each one contains:
/// - A reference to a BLAS, this ***must*** be interacted with using [TlasInstance::new] or [TlasInstance::set_blas], a
/// TlasInstance that references a BLAS keeps that BLAS from being dropped, but if the BLAS is explicitly destroyed (e.g.
/// using [Blas::destroy]) the TlasInstance becomes invalid
/// - A user accessible transformation matrix
/// - A user accessible mask
/// - A user accessible custom index
///
/// [Tlas]: crate::Tlas
/// [TlasPackage]: crate::TlasPackage
#[derive(Debug, Clone)]
pub struct TlasInstance {
pub(crate) blas: Arc<BlasShared>,
/// Affine transform matrix 3x4 (rows x columns, row major order).
pub transform: [f32; 12],
/// Custom index for the instance used inside the shader.
///
/// This must only use the lower 24 bits, if any bits are outside that range (byte 4 does not equal 0) the TlasInstance becomes
/// invalid and generates a validation error when built
pub custom_index: u32,
/// Mask for the instance used inside the shader to filter instances.
/// Reports hit only if `(shader_cull_mask & tlas_instance.mask) != 0u`.
pub mask: u8,
}
impl TlasInstance {
/// Construct TlasInstance.
/// - blas: Reference to the bottom level acceleration structure
/// - transform: Transform buffer offset in bytes (optional, required if transform buffer is present)
/// - custom_index: Custom index for the instance used inside the shader (max 24 bits)
/// - mask: Mask for the instance used inside the shader to filter instances
///
/// Note: while one of these contains a reference to a BLAS that BLAS will not be dropped,
/// but it can still be destroyed. Destroying a BLAS that is referenced by one or more
/// TlasInstance(s) will immediately make them invalid. If one or more of those invalid
/// TlasInstances is inside a TlasPackage that is attempted to be built, the build will
/// generate a validation error.
pub fn new(blas: &Blas, transform: [f32; 12], custom_index: u32, mask: u8) -> Self {
Self {
blas: blas.shared.clone(),
transform,
custom_index,
mask,
}
}
/// Set the bottom level acceleration structure.
///
/// See the note on [TlasInstance] about the
/// guarantees of keeping a BLAS alive.
pub fn set_blas(&mut self, blas: &Blas) {
self.blas = blas.shared.clone();
}
}
pub(crate) struct DynContextTlasInstance<'a> {
pub(crate) blas: &'a Data,
pub(crate) transform: &'a [f32; 12],
pub(crate) custom_index: u32,
pub(crate) mask: u8,
}
/// Context version of [TlasInstance].
#[allow(dead_code)]
pub struct ContextTlasInstance<'a, T: Context> {
pub(crate) blas_data: &'a T::BlasData,
pub(crate) transform: &'a [f32; 12],
pub(crate) custom_index: u32,
pub(crate) mask: u8,
}
#[derive(Debug)]
/// Definition for a triangle geometry for a Bottom Level Acceleration Structure (BLAS).
///
/// The size must match the rest of the structures fields, otherwise the build will fail.
/// (e.g. if index count is present in the size, the index buffer must be present as well.)
pub struct BlasTriangleGeometry<'a> {
/// Sub descriptor for the size defining attributes of a triangle geometry.
pub size: &'a BlasTriangleGeometrySizeDescriptor,
/// Vertex buffer.
pub vertex_buffer: &'a Buffer,
/// Offset into the vertex buffer as a factor of the vertex stride.
pub first_vertex: u32,
/// Vertex stride.
pub vertex_stride: wgt::BufferAddress,
/// Index buffer (optional).
pub index_buffer: Option<&'a Buffer>,
/// Index buffer offset in bytes (optional, required if index buffer is present).
pub index_buffer_offset: Option<wgt::BufferAddress>,
/// Transform buffer containing 3x4 (rows x columns, row major) affine transform matrices `[f32; 12]` (optional).
pub transform_buffer: Option<&'a Buffer>,
/// Transform buffer offset in bytes (optional, required if transform buffer is present).
pub transform_buffer_offset: Option<wgt::BufferAddress>,
}
static_assertions::assert_impl_all!(BlasTriangleGeometry<'_>: WasmNotSendSync);
/// Contains the sets of geometry that go into a [Blas].
pub enum BlasGeometries<'a> {
/// Triangle geometry variant.
TriangleGeometries(Vec<BlasTriangleGeometry<'a>>),
}
static_assertions::assert_impl_all!(BlasGeometries<'_>: WasmNotSendSync);
/// Builds the given sets of geometry into the given [Blas].
pub struct BlasBuildEntry<'a> {
/// Reference to the acceleration structure.
pub blas: &'a Blas,
/// Geometries.
pub geometry: BlasGeometries<'a>,
}
static_assertions::assert_impl_all!(BlasBuildEntry<'_>: WasmNotSendSync);
#[derive(Debug)]
pub(crate) struct BlasShared {
pub(crate) context: Arc<C>,
pub(crate) data: Box<Data>,
}
static_assertions::assert_impl_all!(BlasShared: WasmNotSendSync);
#[derive(Debug)]
/// Bottom Level Acceleration Structure (BLAS).
///
/// A BLAS is a device-specific raytracing acceleration structure that contains geometry data.
///
/// These BLASes are combined with transform in a [TlasInstance] to create a [Tlas].
///
/// [Tlas]: crate::Tlas
pub struct Blas {
pub(crate) handle: Option<u64>,
pub(crate) shared: Arc<BlasShared>,
}
static_assertions::assert_impl_all!(Blas: WasmNotSendSync);
impl Blas {
/// Raw handle to the acceleration structure, used inside raw instance buffers.
pub fn handle(&self) -> Option<u64> {
self.handle
}
/// Destroy the associated native resources as soon as possible.
pub fn destroy(&self) {
DynContext::blas_destroy(&*self.shared.context, self.shared.data.as_ref());
}
}
impl Drop for BlasShared {
fn drop(&mut self) {
if !thread::panicking() {
self.context.blas_drop(self.data.as_ref());
}
}
}
pub(crate) struct DynContextBlasTriangleGeometry<'a> {
pub(crate) size: &'a BlasTriangleGeometrySizeDescriptor,
pub(crate) vertex_buffer: &'a Data,
pub(crate) index_buffer: Option<&'a Data>,
pub(crate) transform_buffer: Option<&'a Data>,
pub(crate) first_vertex: u32,
pub(crate) vertex_stride: wgt::BufferAddress,
pub(crate) index_buffer_offset: Option<wgt::BufferAddress>,
pub(crate) transform_buffer_offset: Option<wgt::BufferAddress>,
}
pub(crate) enum DynContextBlasGeometries<'a> {
TriangleGeometries(Box<dyn Iterator<Item = DynContextBlasTriangleGeometry<'a>> + 'a>),
}
pub(crate) struct DynContextBlasBuildEntry<'a> {
pub(crate) blas_data: &'a Data,
pub(crate) geometries: DynContextBlasGeometries<'a>,
}
/// Context version of [BlasTriangleGeometry].
#[allow(dead_code)]
pub struct ContextBlasTriangleGeometry<'a, T: Context> {
pub(crate) size: &'a BlasTriangleGeometrySizeDescriptor,
pub(crate) vertex_buffer: &'a T::BufferData,
pub(crate) index_buffer: Option<&'a T::BufferData>,
pub(crate) transform_buffer: Option<&'a T::BufferData>,
pub(crate) first_vertex: u32,
pub(crate) vertex_stride: wgt::BufferAddress,
pub(crate) index_buffer_offset: Option<wgt::BufferAddress>,
pub(crate) transform_buffer_offset: Option<wgt::BufferAddress>,
}
/// Context version of [BlasGeometries].
pub enum ContextBlasGeometries<'a, T: Context> {
/// Triangle geometries.
TriangleGeometries(Box<dyn Iterator<Item = ContextBlasTriangleGeometry<'a, T>> + 'a>),
}
/// Context version see [BlasBuildEntry].
#[allow(dead_code)]
pub struct ContextBlasBuildEntry<'a, T: Context> {
pub(crate) blas_data: &'a T::BlasData,
pub(crate) geometries: ContextBlasGeometries<'a, T>,
}

169
wgpu/src/api/command_encoder.rs
View File

@ -55,7 +55,15 @@ pub type ImageCopyTexture<'a> = ImageCopyTextureBase<&'a Texture>;
#[cfg(send_sync)]
static_assertions::assert_impl_all!(ImageCopyTexture<'_>: Send, Sync);
use crate::api::blas::{
BlasBuildEntry, BlasGeometries, BlasTriangleGeometry, DynContextBlasBuildEntry,
DynContextBlasGeometries, DynContextBlasTriangleGeometry, DynContextTlasInstance, TlasInstance,
};
use crate::api::tlas::{
DynContextTlasBuildEntry, DynContextTlasPackage, TlasBuildEntry, TlasPackage,
};
pub use wgt::ImageCopyTextureTagged as ImageCopyTextureTaggedBase;
/// View of a texture which can be used to copy to a texture, including
/// color space and alpha premultiplication information.
///
@ -340,3 +348,164 @@ impl CommandEncoder {
)
}
}
/// [`Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE`] must be enabled on the device in order to call these functions.
impl CommandEncoder {
/// Build bottom and top level acceleration structures.
///
/// Builds the BLASes then the TLASes, but does ***not*** build the BLASes into the TLASes,
/// that must be done by setting a TLAS instance in the TLAS package to one that contains the BLAS (and with an appropriate transform)
///
/// # Validation
///
/// - blas: Iterator of bottom level acceleration structure entries to build.
/// For each entry, the provided size descriptor must be strictly smaller or equal to the descriptor given at BLAS creation, this means:
/// - Less or equal number of geometries
/// - Same kind of geometry (with index buffer or without) (same vertex/index format)
/// - Same flags
/// - Less or equal number of vertices
/// - Less or equal number of indices (if applicable)
/// - tlas: iterator of top level acceleration structure packages to build
/// For each entry:
/// - Each BLAS in each TLAS instance must have been being built in the current call or in a previous call to `build_acceleration_structures` or `build_acceleration_structures_unsafe_tlas`
/// - The number of TLAS instances must be less than or equal to the max number of tlas instances when creating (if creating a package with `TlasPackage::new()` this is already satisfied)
///
/// If the device the command encoder is created from does not have [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE] enabled then a validation error is generated
///
/// A bottom level acceleration structure may be build and used as a reference in a top level acceleration structure in the same invocation of this function.
///
/// # Bind group usage
///
/// When a top level acceleration structure is used in a bind group, some validation takes place:
/// - The top level acceleration structure is valid and has been built.
/// - All the bottom level acceleration structures referenced by the top level acceleration structure are valid and have been built prior,
/// or at same time as the containing top level acceleration structure.
///
/// [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE]: wgt::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
pub fn build_acceleration_structures<'a>(
&mut self,
blas: impl IntoIterator<Item = &'a BlasBuildEntry<'a>>,
tlas: impl IntoIterator<Item = &'a TlasPackage>,
) {
let mut blas = blas.into_iter().map(|e: &BlasBuildEntry<'_>| {
let geometries = match &e.geometry {
BlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries
.iter()
.map(
|tg: &BlasTriangleGeometry<'_>| DynContextBlasTriangleGeometry {
size: tg.size,
vertex_buffer: tg.vertex_buffer.data.as_ref(),
index_buffer: tg
.index_buffer
.map(|index_buffer| index_buffer.data.as_ref()),
transform_buffer: tg
.transform_buffer
.map(|transform_buffer| transform_buffer.data.as_ref()),
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
transform_buffer_offset: tg.transform_buffer_offset,
},
);
DynContextBlasGeometries::TriangleGeometries(Box::new(iter))
}
};
DynContextBlasBuildEntry {
blas_data: e.blas.shared.data.as_ref(),
geometries,
}
});
let mut tlas = tlas.into_iter().map(|e: &TlasPackage| {
let instances = e.instances.iter().map(|instance: &Option<TlasInstance>| {
instance.as_ref().map(|instance| DynContextTlasInstance {
blas: instance.blas.data.as_ref(),
transform: &instance.transform,
custom_index: instance.custom_index,
mask: instance.mask,
})
});
DynContextTlasPackage {
tlas_data: e.tlas.data.as_ref(),
instances: Box::new(instances),
lowest_unmodified: e.lowest_unmodified,
}
});
DynContext::command_encoder_build_acceleration_structures(
&*self.context,
self.data.as_ref(),
&mut blas,
&mut tlas,
);
}
/// Build bottom and top level acceleration structures.
/// See [`CommandEncoder::build_acceleration_structures`] for the safe version and more details. All validation in [`CommandEncoder::build_acceleration_structures`] except that
/// listed under tlas applies here as well.
///
/// # Safety
///
/// - The contents of the raw instance buffer must be valid for the underling api.
/// - All bottom level acceleration structures, referenced in the raw instance buffer must be valid and built,
/// when the corresponding top level acceleration structure is built. (builds may happen in the same invocation of this function).
/// - At the time when the top level acceleration structure is used in a bind group, all associated bottom level acceleration structures must be valid,
/// and built (no later than the time when the top level acceleration structure was built).
pub unsafe fn build_acceleration_structures_unsafe_tlas<'a>(
&mut self,
blas: impl IntoIterator<Item = &'a BlasBuildEntry<'a>>,
tlas: impl IntoIterator<Item = &'a TlasBuildEntry<'a>>,
) {
let mut blas = blas.into_iter().map(|e: &BlasBuildEntry<'_>| {
let geometries = match &e.geometry {
BlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries
.iter()
.map(
|tg: &BlasTriangleGeometry<'_>| DynContextBlasTriangleGeometry {
size: tg.size,
vertex_buffer: tg.vertex_buffer.data.as_ref(),
index_buffer: tg
.index_buffer
.map(|index_buffer| index_buffer.data.as_ref()),
transform_buffer: tg
.transform_buffer
.map(|transform_buffer| transform_buffer.data.as_ref()),
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
transform_buffer_offset: tg.transform_buffer_offset,
},
);
DynContextBlasGeometries::TriangleGeometries(Box::new(iter))
}
};
DynContextBlasBuildEntry {
blas_data: e.blas.shared.data.as_ref(),
geometries,
}
});
let mut tlas = tlas
.into_iter()
.map(|e: &TlasBuildEntry<'_>| DynContextTlasBuildEntry {
tlas_data: e.tlas.data.as_ref(),
instance_buffer_data: e.instance_buffer.data.as_ref(),
instance_count: e.instance_count,
});
DynContext::command_encoder_build_acceleration_structures_unsafe_tlas(
&*self.context,
self.data.as_ref(),
&mut blas,
&mut tlas,
);
}
}

60
wgpu/src/api/device.rs
View File

@ -2,6 +2,8 @@ use std::{error, fmt, future::Future, sync::Arc, thread};
use parking_lot::Mutex;
use crate::api::blas::{Blas, BlasGeometrySizeDescriptors, BlasShared, CreateBlasDescriptor};
use crate::api::tlas::{CreateTlasDescriptor, Tlas};
use crate::context::DynContext;
use crate::*;
@ -514,6 +516,64 @@ impl Device {
}
}
/// [`Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE`] must be enabled on the device in order to call these functions.
impl Device {
/// Create a bottom level acceleration structure, used inside a top level acceleration structure for ray tracing.
/// - `desc`: The descriptor of the acceleration structure.
/// - `sizes`: Size descriptor limiting what can be built into the acceleration structure.
///
/// # Validation
/// If any of the following is not satisfied a validation error is generated
///
/// The device ***must*** have [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE] enabled.
/// if `sizes` is [BlasGeometrySizeDescriptors::Triangles] then the following must be satisfied
/// - For every geometry descriptor (for the purposes this is called `geo_desc`) of `sizes.descriptors` the following must be satisfied:
/// - `geo_desc.vertex_format` must be within allowed formats (allowed formats for a given feature set
/// may be queried with [Features::allowed_vertex_formats_for_blas]).
/// - Both or neither of `geo_desc.index_format` and `geo_desc.index_count` must be provided.
///
/// [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE]: wgt::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
/// [Features::allowed_vertex_formats_for_blas]: wgt::Features::allowed_vertex_formats_for_blas
#[must_use]
pub fn create_blas(
&self,
desc: &CreateBlasDescriptor<'_>,
sizes: BlasGeometrySizeDescriptors,
) -> Blas {
let (handle, data) =
DynContext::device_create_blas(&*self.context, self.data.as_ref(), desc, sizes);
Blas {
#[allow(clippy::arc_with_non_send_sync)]
shared: Arc::new(BlasShared {
context: Arc::clone(&self.context),
data,
}),
handle,
}
}
/// Create a top level acceleration structure, used for ray tracing.
/// - `desc`: The descriptor of the acceleration structure.
///
/// # Validation
/// If any of the following is not satisfied a validation error is generated
///
/// The device ***must*** have [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE] enabled.
///
/// [Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE]: wgt::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
#[must_use]
pub fn create_tlas(&self, desc: &CreateTlasDescriptor<'_>) -> Tlas {
let data = DynContext::device_create_tlas(&*self.context, self.data.as_ref(), desc);
Tlas {
context: Arc::clone(&self.context),
data,
max_instances: desc.max_instances,
}
}
}
impl Drop for Device {
fn drop(&mut self) {
if !thread::panicking() {

4
wgpu/src/api/mod.rs
View File

@ -28,6 +28,7 @@ mod buffer;
mod command_buffer;
mod command_encoder;
// Not a root type, but common descriptor types for pipelines.
mod blas;
mod common_pipeline;
mod compute_pass;
mod compute_pipeline;
@ -47,10 +48,12 @@ mod surface;
mod surface_texture;
mod texture;
mod texture_view;
mod tlas;
pub use adapter::*;
pub use bind_group::*;
pub use bind_group_layout::*;
pub use blas::*;
pub use buffer::*;
pub use command_buffer::*;
pub use command_encoder::*;
@ -73,6 +76,7 @@ pub use surface::*;
pub use surface_texture::*;
pub use texture::*;
pub use texture_view::*;
pub use tlas::*;
/// Object debugging label.
pub type Label<'a> = Option<&'a str>;

198
wgpu/src/api/tlas.rs Normal file
View File

@ -0,0 +1,198 @@
use crate::api::blas::{ContextTlasInstance, DynContextTlasInstance, TlasInstance};
use crate::context::{Context, DynContext};
use crate::{BindingResource, Buffer, Data, Label, C};
use std::ops::{Index, IndexMut, Range};
use std::sync::Arc;
use std::thread;
use wgt::WasmNotSendSync;
/// Descriptor to create top level acceleration structures.
pub type CreateTlasDescriptor<'a> = wgt::CreateTlasDescriptor<Label<'a>>;
static_assertions::assert_impl_all!(CreateTlasDescriptor<'_>: Send, Sync);
#[derive(Debug)]
/// Top Level Acceleration Structure (TLAS).
///
/// A TLAS contains a series of [TLAS instances], which are a reference to
/// a BLAS and a transformation matrix placing the geometry in the world.
///
/// A TLAS contains TLAS instances in a device readable form, you cant interact
/// directly with these, instead you have to build the TLAS with [TLAS instances].
///
/// [TLAS instances]: TlasInstance
pub struct Tlas {
pub(crate) context: Arc<C>,
pub(crate) data: Box<Data>,
pub(crate) max_instances: u32,
}
static_assertions::assert_impl_all!(Tlas: WasmNotSendSync);
impl Tlas {
/// Destroy the associated native resources as soon as possible.
pub fn destroy(&self) {
DynContext::tlas_destroy(&*self.context, self.data.as_ref());
}
}
impl Drop for Tlas {
fn drop(&mut self) {
if !thread::panicking() {
self.context.tlas_drop(self.data.as_ref());
}
}
}
/// Entry for a top level acceleration structure build.
/// Used with raw instance buffers for an unvalidated builds.
/// See [TlasPackage] for the safe version.
pub struct TlasBuildEntry<'a> {
/// Reference to the acceleration structure.
pub tlas: &'a Tlas,
/// Reference to the raw instance buffer, each instance is similar to [TlasInstance] but contains a handle to the BLAS.
pub instance_buffer: &'a Buffer,
/// Number of instances in the instance buffer.
pub instance_count: u32,
}
static_assertions::assert_impl_all!(TlasBuildEntry<'_>: WasmNotSendSync);
/// The safe version of TlasEntry, containing TlasInstances instead of a raw buffer.
pub struct TlasPackage {
pub(crate) tlas: Tlas,
pub(crate) instances: Vec<Option<TlasInstance>>,
pub(crate) lowest_unmodified: u32,
}
static_assertions::assert_impl_all!(TlasPackage: WasmNotSendSync);
impl TlasPackage {
/// Construct [TlasPackage] consuming the [Tlas] (prevents modification of the [Tlas] without using this package).
pub fn new(tlas: Tlas) -> Self {
let max_instances = tlas.max_instances;
Self::new_with_instances(tlas, vec![None; max_instances as usize])
}
/// Construct [TlasPackage] consuming the [Tlas] (prevents modification of the Tlas without using this package).
/// This constructor moves the instances into the package (the number of instances needs to fit into tlas,
/// otherwise when building a validation error will be raised).
pub fn new_with_instances(tlas: Tlas, instances: Vec<Option<TlasInstance>>) -> Self {
Self {
tlas,
lowest_unmodified: instances.len() as u32,
instances,
}
}
/// Get a reference to all instances.
pub fn get(&self) -> &[Option<TlasInstance>] {
&self.instances
}
/// Get a mutable slice to a range of instances.
/// Returns None if the range is out of bounds.
/// All elements from the lowest accessed index up are marked as modified.
// this recommendation is not useful yet, but is likely to be when ability to update arrives or possible optimisations for building get implemented.
/// For best performance it is recommended to prefer access to low elements and modify higher elements as little as possible.
/// This can be done by ordering instances from the most to the least used. It is recommended
/// to use [Self::index_mut] unless the option if out of bounds is required
pub fn get_mut_slice(&mut self, range: Range<usize>) -> Option<&mut [Option<TlasInstance>]> {
if range.end > self.instances.len() {
return None;
}
if range.end as u32 > self.lowest_unmodified {
self.lowest_unmodified = range.end as u32;
}
Some(&mut self.instances[range])
}
/// Get a single mutable reference to an instance.
/// Returns None if the range is out of bounds.
/// All elements from the lowest accessed index up are marked as modified.
// this recommendation is not useful yet, but is likely to be when ability to update arrives or possible optimisations for building get implemented.
/// For best performance it is recommended to prefer access to low elements and modify higher elements as little as possible.
/// This can be done by ordering instances from the most to the least used. It is recommended
/// to use [Self::index_mut] unless the option if out of bounds is required
pub fn get_mut_single(&mut self, index: usize) -> Option<&mut Option<TlasInstance>> {
if index >= self.instances.len() {
return None;
}
if index as u32 + 1 > self.lowest_unmodified {
self.lowest_unmodified = index as u32 + 1;
}
Some(&mut self.instances[index])
}
/// Get the binding resource for the underling acceleration structure, to be used when creating a [BindGroup]
///
/// [BindGroup]: super::BindGroup
pub fn as_binding(&self) -> BindingResource<'_> {
BindingResource::AccelerationStructure(&self.tlas)
}
/// Get a reference to the underling [Tlas].
pub fn tlas(&self) -> &Tlas {
&self.tlas
}
}
impl Index<usize> for TlasPackage {
type Output = Option<TlasInstance>;
fn index(&self, index: usize) -> &Self::Output {
self.instances.index(index)
}
}
impl Index<Range<usize>> for TlasPackage {
type Output = [Option<TlasInstance>];
fn index(&self, index: Range<usize>) -> &Self::Output {
self.instances.index(index)
}
}
impl IndexMut<usize> for TlasPackage {
fn index_mut(&mut self, index: usize) -> &mut Self::Output {
let idx = self.instances.index_mut(index);
if index as u32 + 1 > self.lowest_unmodified {
self.lowest_unmodified = index as u32 + 1;
}
idx
}
}
impl IndexMut<Range<usize>> for TlasPackage {
fn index_mut(&mut self, index: Range<usize>) -> &mut Self::Output {
let idx = self.instances.index_mut(index.clone());
if index.end > self.lowest_unmodified as usize {
self.lowest_unmodified = index.end as u32;
}
idx
}
}
pub(crate) struct DynContextTlasBuildEntry<'a> {
pub(crate) tlas_data: &'a Data,
pub(crate) instance_buffer_data: &'a Data,
pub(crate) instance_count: u32,
}
pub(crate) struct DynContextTlasPackage<'a> {
pub(crate) tlas_data: &'a Data,
pub(crate) instances: Box<dyn Iterator<Item = Option<DynContextTlasInstance<'a>>> + 'a>,
pub(crate) lowest_unmodified: u32,
}
/// Context version see [TlasBuildEntry].
#[allow(dead_code)]
pub struct ContextTlasBuildEntry<'a, T: Context> {
pub(crate) tlas_data: &'a T::TlasData,
pub(crate) instance_buffer_data: &'a T::BufferData,
pub(crate) instance_count: u32,
}
/// Context version see [TlasPackage].
#[allow(dead_code)]
pub struct ContextTlasPackage<'a, T: Context> {
pub(crate) tlas_data: &'a T::TlasData,
pub(crate) instances: Box<dyn Iterator<Item = Option<ContextTlasInstance<'a, T>>> + 'a>,
pub(crate) lowest_unmodified: u32,
}

56
wgpu/src/backend/webgpu.rs
View File

@ -1137,6 +1137,8 @@ impl crate::context::Context for ContextWebGpu {
type RenderBundleEncoderData = Sendable<webgpu_sys::GpuRenderBundleEncoder>;
type RenderBundleData = Sendable<webgpu_sys::GpuRenderBundle>;
type SurfaceData = Sendable<(Canvas, webgpu_sys::GpuCanvasContext)>;
type BlasData = ();
type TlasData = ();
type SurfaceOutputDetail = SurfaceOutputDetail;
type SubmissionIndexData = ();
@ -1758,6 +1760,9 @@ impl crate::context::Context for ContextWebGpu {
crate::BindingResource::TextureViewArray(..) => {
panic!("Web backend does not support BINDING_INDEXING extension")
}
crate::BindingResource::AccelerationStructure(_) => {
unimplemented!("Raytracing not implemented for web")
}
};
webgpu_sys::GpuBindGroupEntry::new(binding.binding, &mapped_resource)
@ -3384,6 +3389,57 @@ impl crate::context::Context for ContextWebGpu {
fn render_pass_end(&self, pass_data: &mut Self::RenderPassData) {
pass_data.0.end();
}
fn device_create_blas(
&self,
_device_data: &Self::DeviceData,
_desc: &crate::CreateBlasDescriptor<'_>,
_sizes: wgt::BlasGeometrySizeDescriptors,
) -> (Option<u64>, Self::BlasData) {
unimplemented!("Raytracing not implemented for web");
}
fn device_create_tlas(
&self,
_device_data: &Self::DeviceData,
_desc: &crate::CreateTlasDescriptor<'_>,
) -> Self::TlasData {
unimplemented!("Raytracing not implemented for web");
}
fn command_encoder_build_acceleration_structures_unsafe_tlas<'a>(
&'a self,
_encoder_data: &Self::CommandEncoderData,
_blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
_tlas: impl Iterator<Item = crate::ContextTlasBuildEntry<'a, Self>>,
) {
unimplemented!("Raytracing not implemented for web");
}
fn command_encoder_build_acceleration_structures<'a>(
&'a self,
_encoder_data: &Self::CommandEncoderData,
_blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
_tlas: impl Iterator<Item = crate::ContextTlasPackage<'a, Self>>,
) {
unimplemented!("Raytracing not implemented for web");
}
fn blas_destroy(&self, _blas_data: &Self::BlasData) {
unimplemented!("Raytracing not implemented for web");
}
fn blas_drop(&self, _blas_data: &Self::BlasData) {
unimplemented!("Raytracing not implemented for web");
}
fn tlas_destroy(&self, _tlas_data: &Self::TlasData) {
unimplemented!("Raytracing not implemented for web");
}
fn tlas_drop(&self, _tlas_data: &Self::TlasData) {
unimplemented!("Raytracing not implemented for web");
}
}
pub(crate) type SurfaceOutputDetail = ();

210
wgpu/src/backend/wgpu_core.rs
View File

@ -504,6 +504,18 @@ pub struct CommandEncoder {
open: bool,
}
#[derive(Debug)]
pub struct Blas {
id: wgc::id::BlasId,
// error_sink: ErrorSink,
}
#[derive(Debug)]
pub struct Tlas {
id: wgc::id::TlasId,
// error_sink: ErrorSink,
}
impl crate::Context for ContextWgpuCore {
type AdapterData = wgc::id::AdapterId;
type DeviceData = Device;
@ -532,6 +544,8 @@ impl crate::Context for ContextWgpuCore {
type SubmissionIndexData = wgc::SubmissionIndex;
type RequestAdapterFuture = Ready<Option<Self::AdapterData>>;
type BlasData = Blas;
type TlasData = Tlas;
#[allow(clippy::type_complexity)]
type RequestDeviceFuture =
@ -983,6 +997,11 @@ impl crate::Context for ContextWgpuCore {
&remaining_arrayed_texture_views[array.len()..];
bm::BindingResource::TextureViewArray(Borrowed(slice))
}
BindingResource::AccelerationStructure(acceleration_structure) => {
bm::BindingResource::AccelerationStructure(
downcast_tlas(acceleration_structure).id,
)
}
},
})
.collect::<Vec<_>>();
@ -2984,6 +3003,194 @@ impl crate::Context for ContextWgpuCore {
);
}
}
fn device_create_blas(
&self,
device_data: &Self::DeviceData,
desc: &crate::CreateBlasDescriptor<'_>,
sizes: wgt::BlasGeometrySizeDescriptors,
) -> (Option<u64>, Self::BlasData) {
let global = &self.0;
let (id, handle, error) = global.device_create_blas(
device_data.id,
&desc.map_label(|l| l.map(Borrowed)),
sizes,
None,
);
if let Some(cause) = error {
self.handle_error(
&device_data.error_sink,
cause,
desc.label,
"Device::create_blas",
);
}
(
handle,
Blas {
id,
// error_sink: Arc::clone(&device_data.error_sink),
},
)
}
fn device_create_tlas(
&self,
device_data: &Self::DeviceData,
desc: &crate::CreateTlasDescriptor<'_>,
) -> Self::TlasData {
let global = &self.0;
let (id, error) =
global.device_create_tlas(device_data.id, &desc.map_label(|l| l.map(Borrowed)), None);
if let Some(cause) = error {
self.handle_error(
&device_data.error_sink,
cause,
desc.label,
"Device::create_blas",
);
}
Tlas {
id,
// error_sink: Arc::clone(&device_data.error_sink),
}
}
fn command_encoder_build_acceleration_structures_unsafe_tlas<'a>(
&'a self,
encoder_data: &Self::CommandEncoderData,
blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
tlas: impl Iterator<Item = crate::ContextTlasBuildEntry<'a, Self>>,
) {
let global = &self.0;
let blas = blas.map(|e: crate::ContextBlasBuildEntry<'_, Self>| {
let geometries = match e.geometries {
crate::ContextBlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries.into_iter().map(|tg| {
wgc::ray_tracing::BlasTriangleGeometry {
vertex_buffer: tg.vertex_buffer.id,
index_buffer: tg.index_buffer.map(|buf| buf.id),
transform_buffer: tg.transform_buffer.map(|buf| buf.id),
size: tg.size,
transform_buffer_offset: tg.transform_buffer_offset,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
}
});
wgc::ray_tracing::BlasGeometries::TriangleGeometries(Box::new(iter))
}
};
wgc::ray_tracing::BlasBuildEntry {
blas_id: e.blas_data.id,
geometries,
}
});
let tlas = tlas
.into_iter()
.map(|e: crate::ContextTlasBuildEntry<'a, ContextWgpuCore>| {
wgc::ray_tracing::TlasBuildEntry {
tlas_id: e.tlas_data.id,
instance_buffer_id: e.instance_buffer_data.id,
instance_count: e.instance_count,
}
});
if let Err(cause) = global.command_encoder_build_acceleration_structures_unsafe_tlas(
encoder_data.id,
blas,
tlas,
) {
self.handle_error_nolabel(
&encoder_data.error_sink,
cause,
"CommandEncoder::build_acceleration_structures_unsafe_tlas",
);
}
}
fn command_encoder_build_acceleration_structures<'a>(
&'a self,
encoder_data: &Self::CommandEncoderData,
blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
tlas: impl Iterator<Item = crate::ContextTlasPackage<'a, Self>>,
) {
let global = &self.0;
let blas = blas.map(|e: crate::ContextBlasBuildEntry<'_, Self>| {
let geometries = match e.geometries {
crate::ContextBlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries.into_iter().map(|tg| {
wgc::ray_tracing::BlasTriangleGeometry {
vertex_buffer: tg.vertex_buffer.id,
index_buffer: tg.index_buffer.map(|buf| buf.id),
transform_buffer: tg.transform_buffer.map(|buf| buf.id),
size: tg.size,
transform_buffer_offset: tg.transform_buffer_offset,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
}
});
wgc::ray_tracing::BlasGeometries::TriangleGeometries(Box::new(iter))
}
};
wgc::ray_tracing::BlasBuildEntry {
blas_id: e.blas_data.id,
geometries,
}
});
let tlas = tlas.into_iter().map(|e| {
let instances =
e.instances
.map(|instance: Option<crate::ContextTlasInstance<'_, _>>| {
instance.map(|instance| wgc::ray_tracing::TlasInstance {
blas_id: instance.blas_data.id,
transform: instance.transform,
custom_index: instance.custom_index,
mask: instance.mask,
})
});
wgc::ray_tracing::TlasPackage {
tlas_id: e.tlas_data.id,
instances: Box::new(instances),
lowest_unmodified: e.lowest_unmodified,
}
});
if let Err(cause) =
global.command_encoder_build_acceleration_structures(encoder_data.id, blas, tlas)
{
self.handle_error_nolabel(
&encoder_data.error_sink,
cause,
"CommandEncoder::build_acceleration_structures_unsafe_tlas",
);
}
}
fn blas_destroy(&self, blas_data: &Self::BlasData) {
let global = &self.0;
let _ = global.blas_destroy(blas_data.id);
}
fn blas_drop(&self, blas_data: &Self::BlasData) {
let global = &self.0;
global.blas_drop(blas_data.id)
}
fn tlas_destroy(&self, tlas_data: &Self::TlasData) {
let global = &self.0;
let _ = global.tlas_destroy(tlas_data.id);
}
fn tlas_drop(&self, tlas_data: &Self::TlasData) {
let global = &self.0;
global.tlas_drop(tlas_data.id)
}
}
#[derive(Debug)]
@ -3145,6 +3352,9 @@ fn downcast_texture_view(
) -> &<ContextWgpuCore as crate::Context>::TextureViewData {
downcast_ref(texture_view.data.as_ref())
}
fn downcast_tlas(tlas: &crate::Tlas) -> &<ContextWgpuCore as crate::Context>::TlasData {
downcast_ref(tlas.data.as_ref())
}
fn downcast_sampler(sampler: &crate::Sampler) -> &<ContextWgpuCore as crate::Context>::SamplerData {
downcast_ref(sampler.data.as_ref())
}

204
wgpu/src/context.rs
View File

@ -49,6 +49,9 @@ pub trait Context: Debug + WasmNotSendSync + Sized {
type RenderBundleData: ContextData;
type SurfaceData: ContextData;
type BlasData: ContextData;
type TlasData: ContextData;
type SurfaceOutputDetail: WasmNotSendSync + 'static;
type SubmissionIndexData: ContextData + Copy;
@ -695,6 +698,34 @@ pub trait Context: Debug + WasmNotSendSync + Sized {
render_bundles: &mut dyn Iterator<Item = &Self::RenderBundleData>,
);
fn render_pass_end(&self, pass_data: &mut Self::RenderPassData);
fn device_create_blas(
&self,
device_data: &Self::DeviceData,
desc: &crate::CreateBlasDescriptor<'_>,
sizes: wgt::BlasGeometrySizeDescriptors,
) -> (Option<u64>, Self::BlasData);
fn device_create_tlas(
&self,
device_data: &Self::DeviceData,
desc: &crate::CreateTlasDescriptor<'_>,
) -> Self::TlasData;
fn command_encoder_build_acceleration_structures_unsafe_tlas<'a>(
&'a self,
encoder_data: &Self::CommandEncoderData,
blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
tlas: impl Iterator<Item = crate::ContextTlasBuildEntry<'a, Self>>,
);
fn command_encoder_build_acceleration_structures<'a>(
&'a self,
encoder_data: &Self::CommandEncoderData,
blas: impl Iterator<Item = crate::ContextBlasBuildEntry<'a, Self>>,
tlas: impl Iterator<Item = crate::ContextTlasPackage<'a, Self>>,
);
fn blas_destroy(&self, blas_data: &Self::BlasData);
fn blas_drop(&self, blas_data: &Self::BlasData);
fn tlas_destroy(&self, tlas_data: &Self::TlasData);
fn tlas_drop(&self, tlas_data: &Self::TlasData);
}
pub(crate) fn downcast_ref<T: Debug + WasmNotSendSync + 'static>(data: &crate::Data) -> &T {
@ -1341,6 +1372,33 @@ pub(crate) trait DynContext: Debug + WasmNotSendSync {
pass_data: &mut crate::Data,
render_bundles: &mut dyn Iterator<Item = &crate::Data>,
);
fn device_create_blas(
&self,
device_data: &crate::Data,
desc: &crate::CreateBlasDescriptor<'_>,
sizes: wgt::BlasGeometrySizeDescriptors,
) -> (Option<u64>, Box<crate::Data>);
fn device_create_tlas(
&self,
device_data: &crate::Data,
desc: &crate::CreateTlasDescriptor<'_>,
) -> Box<crate::Data>;
fn command_encoder_build_acceleration_structures_unsafe_tlas(
&self,
encoder_data: &crate::Data,
blas: &mut dyn Iterator<Item = crate::DynContextBlasBuildEntry<'_>>,
tlas: &mut dyn Iterator<Item = crate::DynContextTlasBuildEntry<'_>>,
);
fn command_encoder_build_acceleration_structures(
&self,
encoder_data: &crate::Data,
blas: &mut dyn Iterator<Item = crate::DynContextBlasBuildEntry<'_>>,
tlas: &mut dyn Iterator<Item = crate::DynContextTlasPackage<'_>>,
);
fn blas_destroy(&self, blas_data: &crate::Data);
fn blas_drop(&self, blas_data: &crate::Data);
fn tlas_destroy(&self, tlas_data: &crate::Data);
fn tlas_drop(&self, tlas_data: &crate::Data);
fn render_pass_end(&self, pass_data: &mut crate::Data);
}
@ -2681,6 +2739,152 @@ where
let pass_data = downcast_mut(pass_data);
Context::render_pass_end(self, pass_data)
}
fn device_create_blas(
&self,
device_data: &crate::Data,
desc: &crate::CreateBlasDescriptor<'_>,
sizes: wgt::BlasGeometrySizeDescriptors,
) -> (Option<u64>, Box<crate::Data>) {
let device_data = downcast_ref(device_data);
let (handle, data) = Context::device_create_blas(self, device_data, desc, sizes);
(handle, Box::new(data) as _)
}
fn device_create_tlas(
&self,
device_data: &crate::Data,
desc: &crate::CreateTlasDescriptor<'_>,
) -> Box<crate::Data> {
let device_data = downcast_ref(device_data);
let data = Context::device_create_tlas(self, device_data, desc);
Box::new(data) as _
}
fn command_encoder_build_acceleration_structures_unsafe_tlas(
&self,
encoder_data: &crate::Data,
blas: &mut dyn Iterator<Item = crate::DynContextBlasBuildEntry<'_>>,
tlas: &mut dyn Iterator<Item = crate::DynContextTlasBuildEntry<'_>>,
) {
let encoder_data = downcast_ref(encoder_data);
let blas = blas.into_iter().map(|e| {
let geometries = match e.geometries {
crate::DynContextBlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries.into_iter().map(|tg| {
crate::ContextBlasTriangleGeometry {
vertex_buffer: downcast_ref(tg.vertex_buffer),
index_buffer: tg.index_buffer.map(downcast_ref),
transform_buffer: tg.transform_buffer.map(downcast_ref),
size: tg.size,
transform_buffer_offset: tg.transform_buffer_offset,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
}
});
crate::ContextBlasGeometries::TriangleGeometries(Box::new(iter))
}
};
crate::ContextBlasBuildEntry {
blas_data: downcast_ref(e.blas_data),
// blas_data: downcast_ref(e.blas_data),
geometries,
}
});
let tlas = tlas
.into_iter()
.map(
|e: crate::DynContextTlasBuildEntry<'_>| crate::ContextTlasBuildEntry {
tlas_data: downcast_ref(e.tlas_data),
instance_buffer_data: downcast_ref(e.instance_buffer_data),
instance_count: e.instance_count,
},
);
Context::command_encoder_build_acceleration_structures_unsafe_tlas(
self,
encoder_data,
blas,
tlas,
)
}
fn command_encoder_build_acceleration_structures(
&self,
encoder_data: &crate::Data,
blas: &mut dyn Iterator<Item = crate::DynContextBlasBuildEntry<'_>>,
tlas: &mut dyn Iterator<Item = crate::DynContextTlasPackage<'_>>,
) {
let encoder_data = downcast_ref(encoder_data);
let blas = blas.into_iter().map(|e| {
let geometries = match e.geometries {
crate::DynContextBlasGeometries::TriangleGeometries(triangle_geometries) => {
let iter = triangle_geometries.into_iter().map(|tg| {
crate::ContextBlasTriangleGeometry {
vertex_buffer: downcast_ref(tg.vertex_buffer),
index_buffer: tg.index_buffer.map(downcast_ref),
transform_buffer: tg.transform_buffer.map(downcast_ref),
size: tg.size,
transform_buffer_offset: tg.transform_buffer_offset,
first_vertex: tg.first_vertex,
vertex_stride: tg.vertex_stride,
index_buffer_offset: tg.index_buffer_offset,
}
});
crate::ContextBlasGeometries::TriangleGeometries(Box::new(iter))
}
};
crate::ContextBlasBuildEntry {
blas_data: downcast_ref(e.blas_data),
// blas_data: downcast_ref(e.blas_data),
geometries,
}
});
let tlas = tlas.into_iter().map(|e: crate::DynContextTlasPackage<'_>| {
let instances =
e.instances
.map(|instance: Option<crate::DynContextTlasInstance<'_>>| {
instance.map(|instance| crate::ContextTlasInstance {
blas_data: downcast_ref(instance.blas),
transform: instance.transform,
custom_index: instance.custom_index,
mask: instance.mask,
})
});
crate::ContextTlasPackage {
tlas_data: downcast_ref(e.tlas_data),
instances: Box::new(instances),
lowest_unmodified: e.lowest_unmodified,
}
});
Context::command_encoder_build_acceleration_structures(self, encoder_data, blas, tlas)
}
fn blas_destroy(&self, blas_data: &crate::Data) {
let blas_data = downcast_ref(blas_data);
Context::blas_destroy(self, blas_data)
}
fn blas_drop(&self, blas_data: &crate::Data) {
let blas_data = downcast_ref(blas_data);
Context::blas_drop(self, blas_data)
}
fn tlas_destroy(&self, tlas_data: &crate::Data) {
let tlas_data = downcast_ref(tlas_data);
Context::tlas_destroy(self, tlas_data)
}
fn tlas_drop(&self, tlas_data: &crate::Data) {
let tlas_data = downcast_ref(tlas_data);
Context::tlas_drop(self, tlas_data)
}
}
pub trait QueueWriteBuffer: WasmNotSendSync + Debug {