Restore nannou_wgpu and wgpu examples.

Cargo.toml

@@ -14,6 +14,7 @@ members = [
     "nannou_new",
     "nannou_osc",
     "nannou_package",
+    "nannou_wgpu",
     "nature_of_code",
     "scripts/run_all_examples",
     "scripts/set_version",
@@ -26,9 +27,11 @@ resolver = "2"
 bevy = "0.14.0"
 bevy_egui = "0.28.0"
 bevy-inspector-egui = "0.25.0"
+image = "0.25"
 rayon = "1.10"
 bevy_common_assets = "0.11.0"
 serde = "1"
 serde_json = "1"
 toml = "0.8"
-serde_yaml = "0.9"
+serde_yaml = "0.9"
+wgpu = "0.20"

examples/Cargo.toml

@@ -230,4 +230,24 @@ required-features = ["video"]
 [[example]]
 name = "video_material"
 path = "video/video_material.rs"
-required-features = ["video"]
+required-features = ["video"]
+
+# WebGPU
+[[example]]
+name = "wgpu_compute_shader"
+path = "wgpu/wgpu_compute_shader/wgpu_compute_shader.rs"
+[[example]]
+name = "wgpu_image"
+path = "wgpu/wgpu_image/wgpu_image.rs"
+[[example]]
+name = "wgpu_image_sequence"
+path = "wgpu/wgpu_image_sequence/wgpu_image_sequence.rs"
+[[example]]
+name = "wgpu_instancing"
+path = "wgpu/wgpu_instancing/wgpu_instancing.rs"
+[[example]]
+name = "wgpu_teapot"
+path = "wgpu/wgpu_teapot/wgpu_teapot.rs"
+[[example]]
+name = "wgpu_triangle"
+path = "wgpu/wgpu_triangle/wgpu_triangle.rs"

examples/wgpu/wgpu_compute_shader/shaders/cs.wgsl

@@ -0,0 +1,22 @@
+struct Buffer {
+    data: array<f32>,
+};
+
+struct Uniforms {
+    time: f32,
+    freq: f32,
+    oscillator_count: u32,
+};
+
+@group(0) @binding(0)
+var<storage, read_write> output: Buffer;
+@group(0) @binding(1)
+var<uniform> uniforms: Uniforms;
+
+@compute @workgroup_size(1, 1, 1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    let index: u32 = id.x;
+    let phase: f32 = uniforms.time + f32(index) * uniforms.freq / f32(uniforms.oscillator_count);
+    output.data[index] = sin(phase) * 0.5 + 0.5;
+    return;
+}

examples/wgpu/wgpu_compute_shader/wgpu_compute_shader.rs

@@ -0,0 +1,282 @@
+//! A small GPU compute shader demonstration.
+//!
+//! Here we use a compute shader to calculate the amplitude of `OSCILLATOR_COUNT` number of
+//! oscillators. The oscillator amplitudes are then laid out across the screen using rectangles
+//! with a gray value equal to the amplitude. Real-time interaction is demonstrated by providing
+//! access to time, frequency (mouse `x`) and the number of oscillators via uniform data.
+
+use nannou::prelude::*;
+use nannou::wgpu::BufferInitDescriptor;
+use std::sync::{Arc, Mutex};
+
+struct Model {
+    compute: Compute,
+    oscillators: Arc<Mutex<Vec<f32>>>,
+}
+
+struct Compute {
+    oscillator_buffer: wgpu::Buffer,
+    oscillator_buffer_size: wgpu::BufferAddress,
+    uniform_buffer: wgpu::Buffer,
+    bind_group: wgpu::BindGroup,
+    pipeline: wgpu::ComputePipeline,
+}
+
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct Uniforms {
+    time: f32,
+    freq: f32,
+    oscillator_count: u32,
+}
+
+const OSCILLATOR_COUNT: u32 = 128;
+
+fn main() {
+    nannou::app(model).update(update).run();
+}
+
+fn model(app: &App) -> Model {
+    let w_id = app.new_window().size(1440, 512).view(view).build().unwrap();
+    let window = app.window(w_id).unwrap();
+    let device = window.device();
+
+    // Create the compute shader module.
+    let cs_desc = wgpu::include_wgsl!("shaders/cs.wgsl");
+    let cs_mod = device.create_shader_module(cs_desc);
+
+    // Create the buffer that will store the result of our compute operation.
+    let oscillator_buffer_size =
+        (OSCILLATOR_COUNT as usize * std::mem::size_of::<f32>()) as wgpu::BufferAddress;
+    let oscillator_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("oscillators"),
+        size: oscillator_buffer_size,
+        usage: wgpu::BufferUsages::STORAGE
+            | wgpu::BufferUsages::COPY_DST
+            | wgpu::BufferUsages::COPY_SRC,
+        mapped_at_creation: false,
+    });
+
+    // Create the buffer that will store time.
+    let uniforms = create_uniforms(app.time, app.mouse.x, window.rect());
+    let uniforms_bytes = uniforms_as_bytes(&uniforms);
+    let usage = wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST;
+    let uniform_buffer = device.create_buffer_init(&BufferInitDescriptor {
+        label: Some("uniform-buffer"),
+        contents: uniforms_bytes,
+        usage,
+    });
+
+    // Create the bind group and pipeline.
+    let bind_group_layout = create_bind_group_layout(device);
+    let bind_group = create_bind_group(
+        device,
+        &bind_group_layout,
+        &oscillator_buffer,
+        oscillator_buffer_size,
+        &uniform_buffer,
+    );
+    let pipeline_layout = create_pipeline_layout(device, &bind_group_layout);
+    let pipeline = create_compute_pipeline(device, &pipeline_layout, &cs_mod);
+
+    let compute = Compute {
+        oscillator_buffer,
+        oscillator_buffer_size,
+        uniform_buffer,
+        bind_group,
+        pipeline,
+    };
+
+    // The vector that we will write oscillator values to.
+    let oscillators = Arc::new(Mutex::new(vec![0.0; OSCILLATOR_COUNT as usize]));
+
+    Model {
+        compute,
+        oscillators,
+    }
+}
+
+fn update(app: &App, model: &mut Model, _update: Update) {
+    let window = app.main_window();
+    let device = window.device();
+    let win_rect = window.rect();
+    let compute = &mut model.compute;
+
+    // The buffer into which we'll read some data.
+    let read_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("read-oscillators"),
+        size: compute.oscillator_buffer_size,
+        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+        mapped_at_creation: false,
+    });
+
+    // An update for the uniform buffer with the current time.
+    let uniforms = create_uniforms(app.time, app.mouse.x, win_rect);
+    let uniforms_size = std::mem::size_of::<Uniforms>() as wgpu::BufferAddress;
+    let uniforms_bytes = uniforms_as_bytes(&uniforms);
+    let usage = wgpu::BufferUsages::COPY_SRC;
+    let new_uniform_buffer = device.create_buffer_init(&BufferInitDescriptor {
+        label: Some("uniform-data-transfer"),
+        contents: uniforms_bytes,
+        usage,
+    });
+
+    // The encoder we'll use to encode the compute pass.
+    let desc = wgpu::CommandEncoderDescriptor {
+        label: Some("oscillator-compute"),
+    };
+    let mut encoder = device.create_command_encoder(&desc);
+    encoder.copy_buffer_to_buffer(
+        &new_uniform_buffer,
+        0,
+        &compute.uniform_buffer,
+        0,
+        uniforms_size,
+    );
+    {
+        let pass_desc = wgpu::ComputePassDescriptor {
+            label: Some("nannou-wgpu_compute_shader-compute_pass"),
+        };
+        let mut cpass = encoder.begin_compute_pass(&pass_desc);
+        cpass.set_pipeline(&compute.pipeline);
+        cpass.set_bind_group(0, &compute.bind_group, &[]);
+        cpass.dispatch_workgroups(OSCILLATOR_COUNT as u32, 1, 1);
+    }
+    encoder.copy_buffer_to_buffer(
+        &compute.oscillator_buffer,
+        0,
+        &read_buffer,
+        0,
+        compute.oscillator_buffer_size,
+    );
+
+    // Submit the compute pass to the device's queue.
+    window.queue().submit(Some(encoder.finish()));
+
+    // Spawn a future that reads the result of the compute pass.
+    let oscillators = model.oscillators.clone();
+    let future = async move {
+        let slice = read_buffer.slice(..);
+        let (tx, rx) = futures::channel::oneshot::channel();
+        slice.map_async(wgpu::MapMode::Read, |res| {
+            tx.send(res).expect("The channel was closed");
+        });
+        if let Ok(_) = rx.await {
+            if let Ok(mut oscillators) = oscillators.lock() {
+                let bytes = &slice.get_mapped_range()[..];
+                // "Cast" the slice of bytes to a slice of floats as required.
+                let floats = {
+                    let len = bytes.len() / std::mem::size_of::<f32>();
+                    let ptr = bytes.as_ptr() as *const f32;
+                    unsafe { std::slice::from_raw_parts(ptr, len) }
+                };
+                oscillators.copy_from_slice(floats);
+            }
+        }
+    };
+    tokio::spawn(future);
+
+    // Check for resource cleanups and mapping callbacks.
+    //
+    // Note that this line is not necessary in our case, as the device we are using already gets
+    // polled when nannou submits the command buffer for drawing and presentation after `view`
+    // completes. If we were to use a standalone device to create our buffer and perform our
+    // compute (rather than the device requested during window creation), calling `poll` regularly
+    // would be a must.
+    //
+    // device.poll(false);
+}
+
+fn view(app: &App, model: &Model, frame: Frame) {
+    frame.clear(BLACK);
+    let draw = app.draw();
+    let window = app.window(frame.window_id()).unwrap();
+    let rect = window.rect();
+
+    if let Ok(oscillators) = model.oscillators.lock() {
+        let w = rect.w() / OSCILLATOR_COUNT as f32;
+        let h = rect.h();
+        let half_w = w * 0.5;
+        for (i, &osc) in oscillators.iter().enumerate() {
+            let x = half_w + map_range(i as u32, 0, OSCILLATOR_COUNT, rect.left(), rect.right());
+            draw.rect().w_h(w, h).x(x).color(gray(osc));
+        }
+    }
+
+    draw.to_frame(app, &frame).unwrap();
+}
+
+fn create_uniforms(time: f32, mouse_x: f32, win_rect: geom::Rect) -> Uniforms {
+    let freq = map_range(
+        mouse_x,
+        win_rect.left(),
+        win_rect.right(),
+        0.0,
+        win_rect.w(),
+    );
+    let oscillator_count = OSCILLATOR_COUNT;
+    Uniforms {
+        time,
+        freq,
+        oscillator_count,
+    }
+}
+
+fn create_bind_group_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {
+    let storage_dynamic = false;
+    let storage_readonly = false;
+    let uniform_dynamic = false;
+    wgpu::BindGroupLayoutBuilder::new()
+        .storage_buffer(
+            wgpu::ShaderStages::COMPUTE,
+            storage_dynamic,
+            storage_readonly,
+        )
+        .uniform_buffer(wgpu::ShaderStages::COMPUTE, uniform_dynamic)
+        .build(device)
+}
+
+fn create_bind_group(
+    device: &wgpu::Device,
+    layout: &wgpu::BindGroupLayout,
+    oscillator_buffer: &wgpu::Buffer,
+    oscillator_buffer_size: wgpu::BufferAddress,
+    uniform_buffer: &wgpu::Buffer,
+) -> wgpu::BindGroup {
+    let buffer_size_bytes = std::num::NonZeroU64::new(oscillator_buffer_size).unwrap();
+    wgpu::BindGroupBuilder::new()
+        .buffer_bytes(oscillator_buffer, 0, Some(buffer_size_bytes))
+        .buffer::<Uniforms>(uniform_buffer, 0..1)
+        .build(device, layout)
+}
+
+fn create_pipeline_layout(
+    device: &wgpu::Device,
+    bind_group_layout: &wgpu::BindGroupLayout,
+) -> wgpu::PipelineLayout {
+    device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+        label: Some("nannou"),
+        bind_group_layouts: &[&bind_group_layout],
+        push_constant_ranges: &[],
+    })
+}
+
+fn create_compute_pipeline(
+    device: &wgpu::Device,
+    layout: &wgpu::PipelineLayout,
+    cs_mod: &wgpu::ShaderModule,
+) -> wgpu::ComputePipeline {
+    let desc = wgpu::ComputePipelineDescriptor {
+        label: Some("nannou"),
+        layout: Some(layout),
+        module: &cs_mod,
+        entry_point: "main",
+    };
+    device.create_compute_pipeline(&desc)
+}
+
+// See `nannou::wgpu::bytes` docs for why these are necessary.
+
+fn uniforms_as_bytes(uniforms: &Uniforms) -> &[u8] {
+    unsafe { wgpu::bytes::from(uniforms) }
+}

examples/wgpu/wgpu_image/shaders/fs.wgsl

@@ -0,0 +1,14 @@
+struct FragmentOutput {
+    @location(0) f_color: vec4<f32>,
+};
+
+@group(0) @binding(0)
+var tex: texture_2d<f32>;
+@group(0) @binding(1)
+var tex_sampler: sampler;
+
+@fragment
+fn main(@location(0) tex_coords: vec2<f32>) -> FragmentOutput {
+    let out_color: vec4<f32> = textureSample(tex, tex_sampler, tex_coords);
+    return FragmentOutput(out_color);
+}

examples/wgpu/wgpu_image/shaders/vs.wgsl

@@ -0,0 +1,11 @@
+struct VertexOutput {
+    @location(0) tex_coords: vec2<f32>,
+    @builtin(position) out_pos: vec4<f32>,
+};
+
+@vertex
+fn main(@location(0) pos: vec2<f32>) -> VertexOutput {
+    let tex_coords: vec2<f32> = vec2<f32>(pos.x * 0.5 + 0.5, 1.0 - (pos.y * 0.5 + 0.5));
+    let out_pos: vec4<f32> = vec4<f32>(pos, 0.0, 1.0);
+    return VertexOutput(tex_coords, out_pos);
+}