Merge pull request #9 from victorvde/main

Update compute-shader-toy and compute-shader hello to wgpu 0.10

Author: raphlinus

compute-shader-hello/Cargo.lock

@@ -4,9 +4,10 @@ version = "0.1.0"
 authors = ["Raph Levien <raph@google.com>"]
 license = "MIT/Apache-2.0"
 edition = "2018"
+resolver = "2"
 
 [dependencies]
-wgpu = "0.8"
+wgpu = "0.11.0"
 env_logger = "0.8"
 pollster = "0.2"
-bytemuck = { version = "1.4", features = ["derive"] }
+bytemuck = { version = "1.7", features = ["derive"] }

compute-shader-hello/Cargo.toml

@@ -16,16 +16,14 @@
 
 //! A simple application to run a compute shader.
 
-mod encode;
-
 use std::time::Instant;
 
 use wgpu::util::DeviceExt;
 
-use encode::Codable;
+use bytemuck;
 
 async fn run() {
-    let instance = wgpu::Instance::new(wgpu::BackendBit::PRIMARY);
+    let instance = wgpu::Instance::new(wgpu::Backends::PRIMARY);
     let adapter = instance.request_adapter(&Default::default()).await.unwrap();
     let features = adapter.features();
     let (device, queue) = adapter
@@ -39,17 +37,11 @@ async fn run() {
         )
         .await
         .unwrap();
-    let mut shader_flags = wgpu::ShaderFlags::VALIDATION;
-    if matches!(
-        adapter.get_info().backend,
-        wgpu::Backend::Vulkan | wgpu::Backend::Metal | wgpu::Backend::Gl
-    ) {
-        shader_flags |= wgpu::ShaderFlags::EXPERIMENTAL_TRANSLATION;
-    }
     let query_set = if features.contains(wgpu::Features::TIMESTAMP_QUERY) {
         Some(device.create_query_set(&wgpu::QuerySetDescriptor {
             count: 2,
             ty: wgpu::QueryType::Timestamp,
+            label: None,
         }))
     } else {
         None
@@ -60,28 +52,28 @@ async fn run() {
         label: None,
         //source: wgpu::ShaderSource::SpirV(bytes_to_u32(include_bytes!("alu.spv")).into()),
         source: wgpu::ShaderSource::Wgsl(include_str!("shader.wgsl").into()),
-        flags: shader_flags,
     });
     println!("shader compilation {:?}", start_instant.elapsed());
-    let input: Vec<u8> = Codable::encode_vec(&[1.0f32, 2.0f32]);
+    let input_f = &[1.0f32, 2.0f32];
+    let input : &[u8] = bytemuck::bytes_of(input_f);
     let input_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
         label: None,
-        contents: &input,
-        usage: wgpu::BufferUsage::STORAGE
-            | wgpu::BufferUsage::COPY_DST
-            | wgpu::BufferUsage::COPY_SRC,
+        contents: input,
+        usage: wgpu::BufferUsages::STORAGE
+            | wgpu::BufferUsages::COPY_DST
+            | wgpu::BufferUsages::COPY_SRC,
     });
     let output_buf = device.create_buffer(&wgpu::BufferDescriptor {
         label: None,
         size: input.len() as u64,
-        usage: wgpu::BufferUsage::MAP_READ | wgpu::BufferUsage::COPY_DST,
+        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
         mapped_at_creation: false,
     });
     // This works if the buffer is initialized, otherwise reads all 0, for some reason.
     let query_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
         label: None,
         contents: &[0; 16],
-        usage: wgpu::BufferUsage::MAP_READ | wgpu::BufferUsage::COPY_DST,
+        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
     });
 
     let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
@@ -109,7 +101,7 @@ async fn run() {
         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(input_f.len() as u32, 1, 1);
     }
     if let Some(query_set) = &query_set {
         encoder.write_timestamp(query_set, 1);
@@ -128,32 +120,18 @@ async fn run() {
     device.poll(wgpu::Maintain::Wait);
     println!("post-poll {:?}", std::time::Instant::now());
     if buf_future.await.is_ok() {
-        let data = buf_slice.get_mapped_range();
+        let data_raw = &*buf_slice.get_mapped_range();
+        let data : &[f32] = bytemuck::cast_slice(data_raw);
         println!("data: {:?}", &*data);
     }
     if features.contains(wgpu::Features::TIMESTAMP_QUERY) {
         let ts_period = queue.get_timestamp_period();
-        let ts_data: Vec<u64> = Codable::decode_vec(&*query_slice.get_mapped_range());
-        let ts_data = ts_data
-            .iter()
-            .map(|ts| *ts as f64 * ts_period as f64 * 1e-6)
-            .collect::<Vec<_>>();
-        println!("compute shader elapsed: {:?}ms", ts_data[1] - ts_data[0]);
+        let ts_data_raw = &*query_slice.get_mapped_range();
+        let ts_data : &[u64] = bytemuck::cast_slice(ts_data_raw);
+        println!("compute shader elapsed: {:?}ms", (ts_data[1] - ts_data[0]) as f64 * ts_period as f64 * 1e-6);
     }
 }
 
-#[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()
-}
-
 fn main() {
     pollster::block_on(run());
 }

compute-shader-hello/src/encode.rs

@@ -20,7 +20,7 @@ struct DataBuf {
 };
 
 [[group(0), binding(0)]]
-var<storage> v_indices: [[access(read_write)]] DataBuf;
+var<storage, read_write> v_indices: DataBuf;
 
 [[stage(compute), workgroup_size(1)]]
 fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {