Add Extended Examples

CHANGELOG.md

@@ -58,6 +58,10 @@ Bottom level categories:
 
 - Empty scissor rects are allowed now, matching the specification. by @PJB3005 in [#3863](https://github.com/gfx-rs/wgpu/pull/3863).
 
+#### General
+
+- Added example repeated-compute. By @JustAnotherCodemonkey in [#3885](https://github.com/gfx-rs/wgpu/pull/3885).
+
 ### Documentation
 
 - Better documentation for draw, draw_indexed, set_viewport and set_scissor_rect. By @genusistimelord in [#3860](https://github.com/gfx-rs/wgpu/pull/3860)

examples/repeated-compute/Cargo.toml

@@ -0,0 +1,29 @@
+[package]
+name = "wgpu-repeated-compute-example"
+version.workspace = true
+license.workspace = true
+edition.workspace = true
+description = "wgpu repeated compute example"
+publish = false
+
+[[bin]]
+name = "repeated-compute"
+path = "src/main.rs"
+
+[dependencies]
+bytemuck.workspace = true
+env_logger.workspace = true
+futures-intrusive.workspace = true
+log.workspace = true
+pollster.workspace = true
+rand = "0.8.5"
+wgpu.workspace = true
+
+[target.'cfg(target_arch = "wasm32")'.dependencies]
+console_error_panic_hook.workspace = true
+console_log.workspace = true
+wasm-bindgen-futures.workspace = true
+
+[dev-dependencies]
+wasm-bindgen-test.workspace = true
+wgpu-test.workspace = true

examples/repeated-compute/README.md

@@ -0,0 +1,4 @@
+# repeated-compute
+
+Repeatedly performs the Collatz calculation used in `hello-compute` on sets of
+random numbers and reports the timings of core events to log.

examples/repeated-compute/src/main.rs

@@ -0,0 +1,238 @@
+//! See hello-compute example main.rs for more details
+//! as similar items here are not explained.
+//!
+//! This example does elaborate on some things though that the
+//! hello-compute example does not such as mapping buffers
+//! and why use the async channels.
+
+use rand::Rng;
+use std::mem::size_of_val;
+
+const OVERFLOW: u32 = 0xffffffff;
+
+async fn run() {
+    let mut numbers = [0u32; 256];
+    let context = WgpuContext::new(size_of_val(&numbers)).await;
+    // Not sure how compatible this is with web but I don't know how to check.
+    let mut rng = rand::thread_rng();
+
+    for _ in 0..10 {
+        for p in numbers.iter_mut() {
+            *p = rng.gen::<u16>() as u32;
+        }
+
+        compute(&mut numbers, &context).await;
+
+        let printed_numbers = numbers
+            .iter()
+            .map(|n| match n {
+                &OVERFLOW => "(overflow)".to_string(),
+                n => n.to_string(),
+            })
+            .collect::<Vec<String>>();
+        log::info!("Results: {printed_numbers:?}");
+    }
+}
+
+async fn compute(local_buffer: &mut [u32], context: &WgpuContext) {
+    log::info!("Beginning GPU compute on data {local_buffer:?}.");
+    // Local buffer contents -> GPU storage buffer
+    /* Adds a write buffer command to the queue. This command is more complicated
+    than it appears. */
+    context.queue.write_buffer(
+        &context.storage_buffer,
+        0,
+        bytemuck::cast_slice(local_buffer),
+    );
+    log::info!("Wrote to buffer.");
+
+    let mut command_encoder = context
+        .device
+        .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+    {
+        let mut compute_pass =
+            command_encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: None });
+        compute_pass.set_pipeline(&context.pipeline);
+        compute_pass.set_bind_group(0, &context.bind_group, &[]);
+        compute_pass.dispatch_workgroups(local_buffer.len() as u32, 1, 1);
+    }
+    // We finish the compute pass by dropping it.
+
+    // Entire storage buffer -> staging buffer.
+    command_encoder.copy_buffer_to_buffer(
+        &context.storage_buffer,
+        0,
+        &context.output_staging_buffer,
+        0,
+        context.storage_buffer.size(),
+    );
+
+    // Finalize the command encoder, add the contained commands to the queue and flush.
+    context.queue.submit(Some(command_encoder.finish()));
+    log::info!("Submitted commands.");
+
+    // Finally time to get our results.
+    /* First we get a buffer slice which represents a chunk of the buffer (which we
+    can't access yet). */
+    // We want the whole thing so use unbounded range.
+    let buffer_slice = context.output_staging_buffer.slice(..);
+    /* Now things get complicated. WebGPU, for safety reasons, only allows either the GPU
+    or CPU to access a buffer's contents at a time. We need to "map" the buffer which means
+    flipping ownership of the buffer over to the CPU and making access legal. We do this
+    with `BufferSlice::map_async`. The problem is that map_async is not an async function
+    and we can't await it. What we need to do instead is pass in a closure that will
+    be executed when the slice is either mapped or the mapping has failed. The problem
+    with this is that we don't have a reliable way to wait in the main code for the buffer
+    to be mapped and even worse, calling get_mapped_range or get_mapped_range_mut prematurely
+    will cause a panic, not return an error. Using async channels solves this as awaiting
+    the recieving of a message from the passed closure will force the outside code to wait.
+    It also doesn't hurt if the closure finishes before the outside code catches up as
+    the message is buffered and recieving will just pick that up. */
+    let (sender, reciever) = futures_intrusive::channel::shared::oneshot_channel();
+    buffer_slice.map_async(wgpu::MapMode::Read, move |r| sender.send(r).unwrap());
+    /* In order for the mapping to be completed, one of three things must happen.
+    One of those can be calling `Device::poll`. This isn't nessecary on the web as devices
+    are polled automatically but natively, we need to make sure this happens manually. */
+    // `Maintain::Wait` will cause the thread to wait on native but not the web.
+    context.device.poll(wgpu::Maintain::Wait);
+    log::info!("Device polled.");
+    // Now we await the recieving and panic if anything went wrong because we're lazy.
+    reciever.receive().await.unwrap().unwrap();
+    log::info!("Result recieved.");
+    // NOW we can call get_mapped_range.
+    {
+        let view = buffer_slice.get_mapped_range();
+        local_buffer.copy_from_slice(bytemuck::cast_slice(&view));
+    }
+    log::info!("Results written to local buffer.");
+    /* We need to make sure all `BufferView`'s are dropped before we do what we're about
+    to do. */
+    // Unmap so that we can copy to the staging buffer in the next iteration.
+    context.output_staging_buffer.unmap();
+}
+
+fn main() {
+    #[cfg(not(target_arch = "wasm32"))]
+    {
+        env_logger::builder()
+            .filter_level(log::LevelFilter::Info)
+            .format_timestamp_nanos()
+            .init();
+        pollster::block_on(run());
+    }
+    #[cfg(target_arch = "wasm32")]
+    {
+        std::panic::set_hook(Box::new(console_error_panic_hook::hook));
+        console_log::init_with_level(log::Level::Info).expect("could not initialize logger");
+        wasm_bindgen_futures::spawn_local(run());
+    }
+}
+
+/// A convenient way to hold together all the useful wgpu stuff together.
+struct WgpuContext {
+    device: wgpu::Device,
+    queue: wgpu::Queue,
+    pipeline: wgpu::ComputePipeline,
+    bind_group: wgpu::BindGroup,
+    storage_buffer: wgpu::Buffer,
+    output_staging_buffer: wgpu::Buffer,
+}
+
+impl WgpuContext {
+    async fn new(buffer_size: usize) -> WgpuContext {
+        let instance = wgpu::Instance::default();
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions::default())
+            .await
+            .unwrap();
+        let (device, queue) = adapter
+            .request_device(
+                &wgpu::DeviceDescriptor {
+                    label: None,
+                    features: wgpu::Features::empty(),
+                    limits: wgpu::Limits::default(),
+                },
+                None,
+            )
+            .await
+            .unwrap();
+
+        // Our shader, kindly compiled with Naga.
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(include_str!(
+                "shader.wgsl"
+            ))),
+        });
+
+        // This is where the GPU will read from and write to.
+        let storage_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: buffer_size as wgpu::BufferAddress,
+            usage: wgpu::BufferUsages::STORAGE
+                | wgpu::BufferUsages::COPY_DST
+                | wgpu::BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+        /* In WebGPU (unlike many graphics libraries), you aren't allowed to map
+        primary buffers (anthing that has any use that isn't COPY_SRC, COPY_DST,
+        MAP_READ, or MAP_WRITE) (as in get pointers to their memory from the CPU).
+        In WebGPU, the idea is that you copy the data into a specialized reading
+        buffer and then read from there. Same would be true for writes if it weren't
+        for the Queue::write_buffer method. */
+        let output_staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: buffer_size as wgpu::BufferAddress,
+            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+            mapped_at_creation: false,
+        });
+
+        // This can be though of as the function signature for our CPU-GPU function.
+        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: None,
+            entries: &[wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: false },
+                    has_dynamic_offset: false,
+                    // Going to have this be None just to be safe.
+                    min_binding_size: None,
+                },
+                count: None,
+            }],
+        });
+        /* This ties actual resources stored in the GPU to our metaphorical function
+        through the binding slots we defined above. */
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &bind_group_layout,
+            entries: &[wgpu::BindGroupEntry {
+                binding: 0,
+                resource: storage_buffer.as_entire_binding(),
+            }],
+        });
+
+        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: None,
+            bind_group_layouts: &[&bind_group_layout],
+            push_constant_ranges: &[],
+        });
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&pipeline_layout),
+            module: &shader,
+            entry_point: "main",
+        });
+
+        WgpuContext {
+            device,
+            queue,
+            pipeline,
+            bind_group,
+            storage_buffer,
+            output_staging_buffer,
+        }
+    }
+}

examples/repeated-compute/src/shader.wgsl

@@ -0,0 +1,38 @@
+@group(0)
+@binding(0)
+var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
+
+// The Collatz Conjecture states that for any integer n:
+// If n is even, n = n/2
+// If n is odd, n = 3n+1
+// And repeat this process for each new n, you will always eventually reach 1.
+// Though the conjecture has not been proven, no counterexample has ever been found.
+// This function returns how many times this recurrence needs to be applied to reach 1.
+fn collatz_iterations(n_base: u32) -> u32{
+    var n: u32 = n_base;
+    var i: u32 = 0u;
+    loop {
+        if (n <= 1u) {
+            break;
+        }
+        if (n % 2u == 0u) {
+            n = n / 2u;
+        }
+        else {
+            // Overflow? (i.e. 3*n + 1 > 0xffffffffu?)
+            if (n >= 1431655765u) {   // 0x55555555u
+                return 4294967295u;   // 0xffffffffu
+            }
+
+            n = 3u * n + 1u;
+        }
+        i = i + 1u;
+    }
+    return i;
+}
+
+@compute
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
+}