Merge remote-tracking branch 'pull_Vecvec/hal-acceleration-structures…

…' into ray-tracing

# Conflicts:
#	Cargo.lock
#	wgpu-hal/Cargo.toml
#	wgpu-hal/src/dx11/mod.rs
#	wgpu-types/src/lib.rs

wgpu-hal/Cargo.toml

@@ -180,10 +180,10 @@ features = ["wgsl-in"]
 [dev-dependencies]
 cfg-if = "1"
 env_logger = "0.10"
+glam = "0.24.2" # for ray-traced-triangle example
 winit = { version = "0.29.4", features = [
     "android-native-activity",
 ] } # for "halmark" example
-glam = "0.21.3" # for ray-traced-triangle example
 
 [target.'cfg(not(target_arch = "wasm32"))'.dev-dependencies]
 glutin = "0.29.1" # for "gles" example

wgpu-hal/examples/ray-traced-triangle/main.rs

@@ -3,16 +3,15 @@ extern crate wgpu_hal as hal;
 use hal::{
     Adapter as _, CommandEncoder as _, Device as _, Instance as _, Queue as _, Surface as _,
 };
-use raw_window_handle::{HasRawDisplayHandle, HasRawWindowHandle};
+use raw_window_handle::{HasDisplayHandle, HasWindowHandle};
 
 use glam::{Affine3A, Mat4, Vec3};
 use std::{
     borrow::{Borrow, Cow},
-    iter, mem,
-    mem::{align_of, size_of},
-    ptr::{self, copy_nonoverlapping},
+    iter, mem, ptr,
     time::Instant,
 };
+use winit::window::WindowButtons;
 
 const COMMAND_BUFFER_PER_CONTEXT: usize = 100;
 const DESIRED_FRAMES: u32 = 3;
@@ -216,38 +215,42 @@ struct Example<A: hal::Api> {
 }
 
 impl<A: hal::Api> Example<A> {
-    fn init(window: &winit::window::Window) -> Result<Self, hal::InstanceError> {
+    fn init(window: &winit::window::Window) -> Result<Self, Box<dyn std::error::Error>> {
         let instance_desc = hal::InstanceDescriptor {
             name: "example",
-            flags: if cfg!(debug_assertions) {
-                hal::InstanceFlags::all()
-            } else {
-                hal::InstanceFlags::empty()
+            flags: wgt::InstanceFlags::default(),
+            dx12_shader_compiler: wgt::Dx12Compiler::Dxc {
+                dxil_path: None,
+                dxc_path: None,
             },
-            dx12_shader_compiler: wgt::Dx12Compiler::Fxc,
             gles_minor_version: wgt::Gles3MinorVersion::default(),
         };
         let instance = unsafe { A::Instance::init(&instance_desc)? };
-        let mut surface = unsafe {
-            instance
-                .create_surface(window.raw_display_handle(), window.raw_window_handle())
-                .unwrap()
+        let surface = {
+            let raw_window_handle = window.window_handle()?.as_raw();
+            let raw_display_handle = window.display_handle()?.as_raw();
+
+            unsafe {
+                instance
+                    .create_surface(raw_display_handle, raw_window_handle)
+                    .unwrap()
+            }
         };
 
         let (adapter, features) = unsafe {
             let mut adapters = instance.enumerate_adapters();
             if adapters.is_empty() {
-                return Err(hal::InstanceError);
+                panic!("No adapters found");
             }
             let exposed = adapters.swap_remove(0);
             dbg!(exposed.features);
             (exposed.adapter, exposed.features)
         };
-        let surface_caps =
-            unsafe { adapter.surface_capabilities(&surface) }.ok_or(hal::InstanceError)?;
+        let surface_caps = unsafe { adapter.surface_capabilities(&surface) }
+            .expect("Surface doesn't support presentation");
         log::info!("Surface caps: {:#?}", surface_caps);
 
-        let hal::OpenDevice { device, mut queue } =
+        let hal::OpenDevice { device, queue } =
             unsafe { adapter.open(features, &wgt::Limits::default()).unwrap() };
 
         let window_size: (u32, u32) = window.inner_size().into();
@@ -320,47 +323,32 @@ impl<A: hal::Api> Example<A> {
 
         let bgl = unsafe { device.create_bind_group_layout(&bgl_desc).unwrap() };
 
-        pub fn make_spirv_raw(data: &[u8]) -> Cow<[u32]> {
-            const MAGIC_NUMBER: u32 = 0x0723_0203;
-            assert_eq!(
-                data.len() % size_of::<u32>(),
-                0,
-                "data size is not a multiple of 4"
-            );
-
-            //If the data happens to be aligned, directly use the byte array,
-            // otherwise copy the byte array in an owned vector and use that instead.
-            let words = if data.as_ptr().align_offset(align_of::<u32>()) == 0 {
-                let (pre, words, post) = unsafe { data.align_to::<u32>() };
-                debug_assert!(pre.is_empty());
-                debug_assert!(post.is_empty());
-                Cow::from(words)
-            } else {
-                let mut words = vec![0u32; data.len() / size_of::<u32>()];
-                unsafe {
-                    copy_nonoverlapping(data.as_ptr(), words.as_mut_ptr() as *mut u8, data.len());
-                }
-                Cow::from(words)
-            };
-
-            assert_eq!(
-                words[0], MAGIC_NUMBER,
-                "wrong magic word {:x}. Make sure you are using a binary SPIRV file.",
-                words[0]
-            );
-
-            words
-        }
-
+        let naga_shader = {
+            let shader_file = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR"))
+                .join("examples")
+                .join("ray-traced-triangle")
+                .join("shader.wgsl");
+            let source = std::fs::read_to_string(shader_file).unwrap();
+            let module = naga::front::wgsl::Frontend::new().parse(&source).unwrap();
+            let info = naga::valid::Validator::new(
+                naga::valid::ValidationFlags::all(),
+                naga::valid::Capabilities::RAY_QUERY,
+            )
+            .validate(&module)
+            .unwrap();
+            hal::NagaShader {
+                module: Cow::Owned(module),
+                info,
+                debug_source: None,
+            }
+        };
+        let shader_desc = hal::ShaderModuleDescriptor {
+            label: None,
+            runtime_checks: false,
+        };
         let shader_module = unsafe {
             device
-                .create_shader_module(
-                    &hal::ShaderModuleDescriptor {
-                        label: None,
-                        runtime_checks: false,
-                    },
-                    hal::ShaderInput::SpirV(&make_spirv_raw(include_bytes!("shader.comp.spv"))),
-                )
+                .create_shader_module(&shader_desc, hal::ShaderInput::Naga(naga_shader))
                 .unwrap()
         };
 
@@ -567,6 +555,7 @@ impl<A: hal::Api> Example<A> {
             dimension: wgt::TextureViewDimension::D2,
             usage: hal::TextureUses::STORAGE_READ_WRITE | hal::TextureUses::COPY_SRC,
             range: wgt::ImageSubresourceRange::default(),
+            plane: None,
         };
         let texture_view = unsafe { device.create_texture_view(&texture, &view_desc).unwrap() };
 
@@ -899,6 +888,7 @@ impl<A: hal::Api> Example<A> {
             dimension: wgt::TextureViewDimension::D2,
             usage: hal::TextureUses::COPY_DST,
             range: wgt::ImageSubresourceRange::default(),
+            plane: None,
         };
         let surface_tex_view = unsafe {
             self.device
@@ -940,7 +930,7 @@ impl<A: hal::Api> Example<A> {
             ctx.encoder.copy_texture_to_texture(
                 &self.texture,
                 hal::TextureUses::COPY_SRC,
-                &surface_tex.borrow(),
+                surface_tex.borrow(),
                 std::iter::once(hal::TextureCopy {
                     src_base: hal::TextureCopyBase {
                         mip_level: 0,
@@ -973,7 +963,7 @@ impl<A: hal::Api> Example<A> {
                 None
             };
             self.queue.submit(&[&cmd_buf], fence_param).unwrap();
-            self.queue.present(&mut self.surface, surface_tex).unwrap();
+            self.queue.present(&self.surface, surface_tex).unwrap();
             ctx.used_cmd_bufs.push(cmd_buf);
             ctx.used_views.push(surface_tex_view);
         };
@@ -1070,53 +1060,54 @@ cfg_if::cfg_if! {
 fn main() {
     env_logger::init();
 
-    let event_loop = winit::event_loop::EventLoop::new();
+    let event_loop = winit::event_loop::EventLoop::new().unwrap();
     let window = winit::window::WindowBuilder::new()
-        .with_title("hal-bunnymark")
+        .with_title("hal-ray-traced-triangle")
         .with_inner_size(winit::dpi::PhysicalSize {
             width: 512,
             height: 512,
         })
         .with_resizable(false)
+        .with_enabled_buttons(WindowButtons::CLOSE)
         .build(&event_loop)
         .unwrap();
 
     let example_result = Example::<Api>::init(&window);
     let mut example = Some(example_result.expect("Selected backend is not supported"));
 
-    event_loop.run(move |event, _, control_flow| {
-        let _ = &window; // force ownership by the closure
-        *control_flow = winit::event_loop::ControlFlow::Poll;
-        match event {
-            winit::event::Event::RedrawEventsCleared => {
-                window.request_redraw();
-            }
-            winit::event::Event::WindowEvent { event, .. } => match event {
-                winit::event::WindowEvent::KeyboardInput {
-                    input:
-                        winit::event::KeyboardInput {
-                            virtual_keycode: Some(winit::event::VirtualKeyCode::Escape),
-                            state: winit::event::ElementState::Pressed,
-                            ..
-                        },
-                    ..
-                }
-                | winit::event::WindowEvent::CloseRequested => {
-                    *control_flow = winit::event_loop::ControlFlow::Exit;
+    event_loop
+        .run(move |event, target| {
+            let _ = &window; // force ownership by the closure
+            target.set_control_flow(winit::event_loop::ControlFlow::Poll);
+            match event {
+                winit::event::Event::WindowEvent { event, .. } => match event {
+                    winit::event::WindowEvent::CloseRequested => {
+                        target.exit();
+                    }
+                    winit::event::WindowEvent::KeyboardInput { event, .. }
+                        if event.physical_key
+                            == winit::keyboard::PhysicalKey::Code(
+                                winit::keyboard::KeyCode::Escape,
+                            ) =>
+                    {
+                        target.exit();
+                    }
+                    winit::event::WindowEvent::RedrawRequested => {
+                        let ex = example.as_mut().unwrap();
+                        ex.render();
+                    }
+                    _ => {
+                        example.as_mut().unwrap().update(event);
+                    }
+                },
+                winit::event::Event::LoopExiting => {
+                    example.take().unwrap().exit();
                 }
-                _ => {
-                    example.as_mut().unwrap().update(event);
+                winit::event::Event::AboutToWait => {
+                    window.request_redraw();
                 }
-            },
-            winit::event::Event::RedrawRequested(_) => {
-                let ex = example.as_mut().unwrap();
-
-                ex.render();
+                _ => {}
             }
-            winit::event::Event::LoopDestroyed => {
-                example.take().unwrap().exit();
-            }
-            _ => {}
-        }
-    });
+        })
+        .unwrap();
 }

wgpu-hal/examples/ray-traced-triangle/shader.wgsl

@@ -0,0 +1,37 @@
+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);
+}

wgpu-hal/src/dx11/mod.rs

@@ -111,10 +111,10 @@ pub struct PipelineLayout {}
 #[derive(Debug)]
 pub struct ShaderModule {}
 #[derive(Debug)]
-pub struct RenderPipeline {}
-#[derive(Debug)]
 pub struct AccelerationStructure {}
 #[derive(Debug)]
+pub struct RenderPipeline {}
+#[derive(Debug)]
 pub struct ComputePipeline {}
 
 impl crate::Surface<Api> for Surface {

wgpu-hal/src/vulkan/adapter.rs

@@ -1403,6 +1403,10 @@ impl super::Adapter {
                 capabilities.push(spv::Capability::StorageImageWriteWithoutFormat);
             }
 
+            if features.contains(wgt::Features::RAY_QUERY) {
+                capabilities.push(spv::Capability::RayQueryKHR);
+            }
+
             let mut flags = spv::WriterFlags::empty();
             flags.set(
                 spv::WriterFlags::DEBUG,