Merge branch 'main' into piet

Author: raphlinus

compute-shader-toy/Cargo.toml

@@ -8,6 +8,7 @@ edition = "2018"
 
 [dependencies]
 wgpu = { version = "0.8.1", features = ["cross"] }
-winit = "0.24"
+winit = "0.25"
 pollster = "0.2"
 async-executor = "1.0"
+bytemuck = "1.5.1"

compute-shader-toy/README.md

@@ -0,0 +1,12 @@
+# Compute shader toy based on wgpu
+
+This is a starting point for running compute shaders and drawing the image output into a window, based on [wgpu]. You should be able to change the shader (paint.wgsl) and run using simply `cargo run`.
+
+If there are syntax errors in the shader, the error message can be pretty cryptic. It might be useful to run [naga] from the commandline to validate the shader code.
+
+The shading language for this example is [WGSL], translated by naga, but it is possible to run wgpu in native mode with SPIR-V shaders as well. At the time of this writing, compute shaders are blocked on [naga#875], but when the fix for that lands, the experience should be better, and that will also open up features like atomics (and possibly subgroup operations) that are not presently supported by the wgpu stack.
+
+[naga]: https://github.com/gfx-rs/naga
+[wgpu]: https://wgpu.rs/
+[WGSL]: https://gpuweb.github.io/gpuweb/wgsl/
+[naga#975]: https://github.com/gfx-rs/naga/issues/875

compute-shader-toy/foo.comp

@@ -1,3 +1,21 @@
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+// Also licensed under MIT license, at your choice.
+
+// A simple vert/frag shader to copy an image to the swapchain.
+
 struct VertexOutput {
     [[location(0)]] tex_coord: vec2<f32>;
     [[builtin(position)]] position: vec4<f32>;

compute-shader-toy/src/copy.wgsl

@@ -16,7 +16,9 @@
 
 //! A simple compute shader example that draws into a window, based on wgpu.
 
-use wgpu::Extent3d;
+use wgpu::util::DeviceExt;
+use wgpu::{BufferUsage, Extent3d};
+
 use winit::{
     event::{Event, WindowEvent},
     event_loop::{ControlFlow, EventLoop},
@@ -27,12 +29,10 @@ async fn run(event_loop: EventLoop<()>, window: Window) {
     let instance = wgpu::Instance::new(wgpu::BackendBit::PRIMARY);
     let surface = unsafe { instance.create_surface(&window) };
     let adapter = instance
-        .request_adapter(
-            &wgpu::RequestAdapterOptions {
-                power_preference: Default::default(),
-                compatible_surface: Some(&surface),
-            },
-        )
+        .request_adapter(&wgpu::RequestAdapterOptions {
+            power_preference: Default::default(),
+            compatible_surface: Some(&surface),
+        })
         .await
         .expect("error finding adapter");
 
@@ -60,31 +60,32 @@ async fn run(event_loop: EventLoop<()>, window: Window) {
         source: wgpu::ShaderSource::Wgsl(include_str!("copy.wgsl").into()),
         flags: shader_flags,
     });
-    let copy_bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-        label: None,
-        entries: &[
-            wgpu::BindGroupLayoutEntry {
-                binding: 0,
-                visibility: wgpu::ShaderStage::FRAGMENT,
-                ty: wgpu::BindingType::Texture {
-                    multisampled: false,
-                    // Should filterable be false if we want nearest-neighbor?
-                    sample_type: wgpu::TextureSampleType::Float { filterable: true },
-                    view_dimension: wgpu::TextureViewDimension::D2,
+    let copy_bind_group_layout =
+        device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: None,
+            entries: &[
+                wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStage::FRAGMENT,
+                    ty: wgpu::BindingType::Texture {
+                        multisampled: false,
+                        // Should filterable be false if we want nearest-neighbor?
+                        sample_type: wgpu::TextureSampleType::Float { filterable: true },
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                    },
+                    count: None,
                 },
-                count: None,
-            },
-            wgpu::BindGroupLayoutEntry {
-                binding: 1,
-                visibility: wgpu::ShaderStage::FRAGMENT,
-                ty: wgpu::BindingType::Sampler {
-                    filtering: false,
-                    comparison: false,
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStage::FRAGMENT,
+                    ty: wgpu::BindingType::Sampler {
+                        filtering: false,
+                        comparison: false,
+                    },
+                    count: None,
                 },
-                count: None,
-            },
-        ]
-    });
+            ],
+        });
     let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
         label: None,
         bind_group_layouts: &[&copy_bind_group_layout],
@@ -107,77 +108,108 @@ async fn run(event_loop: EventLoop<()>, window: Window) {
         depth_stencil: None,
         multisample: wgpu::MultisampleState::default(),
     });
+
+    let img = device.create_texture(&wgpu::TextureDescriptor {
+        label: None,
+        size: Extent3d {
+            width: size.width,
+            height: size.height,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        usage: wgpu::TextureUsage::STORAGE | wgpu::TextureUsage::SAMPLED,
+    });
+    let img_view = img.create_view(&Default::default());
+
+    const CONFIG_SIZE: u64 = 12;
+
+    let config_dev = device.create_buffer(&wgpu::BufferDescriptor {
+        label: None,
+        size: CONFIG_SIZE,
+        usage: BufferUsage::COPY_DST | BufferUsage::STORAGE,
+        mapped_at_creation: false,
+    });
+    let config_resource = config_dev.as_entire_binding();
+
+    let cs_module = device.create_shader_module(&wgpu::ShaderModuleDescriptor {
+        label: None,
+        source: wgpu::ShaderSource::Wgsl(include_str!("paint.wgsl").into()),
+        flags: shader_flags,
+    });
+    let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+        label: None,
+        layout: None,
+        module: &cs_module,
+        entry_point: "main",
+    });
+    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: config_resource,
+            },
+            wgpu::BindGroupEntry {
+                binding: 1,
+                resource: wgpu::BindingResource::TextureView(&img_view),
+            },
+        ],
+    });
+    let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
+        address_mode_u: wgpu::AddressMode::ClampToEdge,
+        address_mode_v: wgpu::AddressMode::ClampToEdge,
+        address_mode_w: wgpu::AddressMode::ClampToEdge,
+        mag_filter: wgpu::FilterMode::Nearest,
+        min_filter: wgpu::FilterMode::Nearest,
+        mipmap_filter: wgpu::FilterMode::Nearest,
+        ..Default::default()
+    });
+    let copy_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: None,
+        layout: &copy_bind_group_layout,
+        entries: &[
+            wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&img_view),
+            },
+            wgpu::BindGroupEntry {
+                binding: 1,
+                resource: wgpu::BindingResource::Sampler(&sampler),
+            },
+        ],
+    });
+    let start_time = std::time::Instant::now();
+
     event_loop.run(move |event, _, control_flow| {
-        //println!("event {:?}", event);
-        *control_flow = ControlFlow::Wait;
+        // TODO: this may be excessive polling. It really should be synchronized with
+        // swapchain presentation, but that's currently underbaked in wgpu.
+        *control_flow = ControlFlow::Poll;
         match event {
             Event::RedrawRequested(_) => {
                 let frame = swap_chain
                     .get_current_frame()
                     .expect("error getting frame from swap chain")
                     .output;
-                let img = device.create_texture(&wgpu::TextureDescriptor {
-                    label: None,
-                    size: Extent3d {
-                        width: size.width,
-                        height: size.height,
-                        depth_or_array_layers: 1,
-                    },
-                    mip_level_count: 1,
-                    sample_count: 1,
-                    dimension: wgpu::TextureDimension::D2,
-                    format: wgpu::TextureFormat::Rgba8Unorm,
-                    usage: wgpu::TextureUsage::STORAGE | wgpu::TextureUsage::SAMPLED,
-                });
-                let img_view = img.create_view(&Default::default());
-                let cs_module = device.create_shader_module(&wgpu::ShaderModuleDescriptor {
-                    label: None,
-                    source: wgpu::ShaderSource::SpirV(bytes_to_u32(include_bytes!("paint.spv")).into()),
-                    //source: wgpu::ShaderSource::Wgsl(include_str!("paint.wgsl").into()),
-                    flags: shader_flags,
-                });
-                let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                    label: None,
-                    layout: None,
-                    module: &cs_module,
-                    entry_point: "main",
-                });
-                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: wgpu::BindingResource::TextureView(&img_view),
-                    }],
-                });
-                let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
-                    address_mode_u: wgpu::AddressMode::ClampToEdge,
-                    address_mode_v: wgpu::AddressMode::ClampToEdge,
-                    address_mode_w: wgpu::AddressMode::ClampToEdge,
-                    mag_filter: wgpu::FilterMode::Nearest,
-                    min_filter: wgpu::FilterMode::Nearest,
-                    mipmap_filter: wgpu::FilterMode::Nearest,
-                    ..Default::default()
-                });
-                                let copy_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+
+                let i_time: f32 = 0.5 + start_time.elapsed().as_micros() as f32 * 1e-6;
+                let config_data = [size.width, size.height, i_time.to_bits()];
+                let config_host = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
                     label: None,
-                    layout: &copy_bind_group_layout,
-                    entries: &[wgpu::BindGroupEntry {
-                        binding: 0,
-                        resource: wgpu::BindingResource::TextureView(&img_view),
-                    },
-                    wgpu::BindGroupEntry {
-                        binding: 1,
-                        resource: wgpu::BindingResource::Sampler(&sampler),
-                    }],
+                    contents: bytemuck::bytes_of(&config_data),
+                    usage: BufferUsage::COPY_SRC,
                 });
                 let mut encoder = device.create_command_encoder(&Default::default());
+                encoder.copy_buffer_to_buffer(&config_host, 0, &config_dev, 0, CONFIG_SIZE);
                 {
                     let mut cpass = encoder.begin_compute_pass(&Default::default());
                     cpass.set_pipeline(&pipeline);
                     cpass.set_bind_group(0, &bind_group, &[]);
-                    cpass.dispatch(1, 1, 1);
+                    cpass.dispatch(size.width / 16, size.height / 16, 1);
                 }
                 {
                     let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
@@ -198,6 +230,9 @@ async fn run(event_loop: EventLoop<()>, window: Window) {
                 }
                 queue.submit(Some(encoder.finish()));
             }
+            Event::MainEventsCleared => {
+                window.request_redraw();
+            }
             Event::WindowEvent {
                 event: WindowEvent::CloseRequested,
                 ..
@@ -212,16 +247,3 @@ fn main() {
     let window = Window::new(&event_loop).unwrap();
     pollster::block_on(run(event_loop, window));
 }
-
-// This will be used if we have spv shaders.
-#[allow(unused)]
-fn bytes_to_u32(bytes: &[u8]) -> Vec<u32> {
-    bytes
-        .chunks_exact(4)
-        .map(|b| {
-            let mut bytes = [0; 4];
-            bytes.copy_from_slice(b);
-            u32::from_le_bytes(bytes)
-        })
-        .collect()
-}

compute-shader-toy/src/main.rs

@@ -14,11 +14,23 @@
 //
 // Also licensed under MIT license, at your choice.
 
-[[group(0), binding(0)]] var outputTex: [[access(write)]] texture_storage_2d<rgba8unorm>;
+[[block]]
+struct Params {
+    width: u32;
+    height: u32;
+    iTime: f32;
+};
+
+[[group(0), binding(0)]] var params: [[access(read)]] Params;
+[[group(0), binding(1)]] var outputTex: [[access(write)]] texture_storage_2d<rgba8unorm>;
 
 [[stage(compute), workgroup_size(16, 16)]]
 fn main([[builtin(global_invocation_id)]] global_ix: vec3<u32>) {
-    let rgba: vec4<f32> = vec4<f32>(1.0, 0.0, 0.0, 1.0);
-    let write_ix = vec2<i32>(i32(global_ix.x), i32(global_ix.y));
-    textureStore(outputTex, write_ix, rgba);
+    let fragCoord: vec2<f32> = vec2<f32>(global_ix.xy) / vec2<f32>(f32(params.width), f32(params.height))
+        - vec2<f32>(0.5, 0.5);
+
+    // Shadertoy-like code can go here.
+    let fragColor: vec4<f32> = vec4<f32>(fragCoord.x + 0.5, fragCoord.y + 0.5, sin(params.iTime), 1.0);
+
+    textureStore(outputTex, vec2<i32>(global_ix.xy), fragColor);
 }