Update Naga with new storage classes API

player/tests/data/buffer-zero-init-for-binding.wgsl

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

wgpu-core/Cargo.toml

@@ -36,7 +36,7 @@ thiserror = "1"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a7ac13a"
+rev = "300039e"
 features = ["wgsl-in"]
 
 [dependencies.wgt]

wgpu-core/src/validation.rs

@@ -325,7 +325,14 @@ impl Resource {
                             wgt::BufferBindingType::Storage { read_only } => {
                                 let mut global_use = GlobalUse::READ | GlobalUse::QUERY;
                                 global_use.set(GlobalUse::WRITE, !read_only);
-                                (naga::StorageClass::Storage, global_use)
+                                let mut naga_access = naga::StorageAccess::LOAD;
+                                naga_access.set(naga::StorageAccess::STORE, !read_only);
+                                (
+                                    naga::StorageClass::Storage {
+                                        access: naga_access,
+                                    },
+                                    global_use,
+                                )
                             }
                         };
                         if self.class != class {
@@ -432,16 +439,26 @@ impl Resource {
                     } => {
                         let naga_format = map_storage_format_to_naga(format)
                             .ok_or(BindingError::BadStorageFormat(format))?;
-                        let usage = match access {
-                            wgt::StorageTextureAccess::ReadOnly => {
-                                GlobalUse::READ | GlobalUse::QUERY
+                        let (naga_access, usage) = match access {
+                            wgt::StorageTextureAccess::ReadOnly => (
+                                naga::StorageAccess::LOAD,
+                                GlobalUse::READ | GlobalUse::QUERY,
+                            ),
+                            wgt::StorageTextureAccess::WriteOnly => (
+                                naga::StorageAccess::STORE,
+                                GlobalUse::WRITE | GlobalUse::QUERY,
+                            ),
+                            wgt::StorageTextureAccess::ReadWrite => {
+                                (naga::StorageAccess::all(), GlobalUse::all())
                             }
-                            wgt::StorageTextureAccess::WriteOnly => {
-                                GlobalUse::WRITE | GlobalUse::QUERY
-                            }
-                            wgt::StorageTextureAccess::ReadWrite => GlobalUse::all(),
                         };
-                        (naga::ImageClass::Storage(naga_format), usage)
+                        (
+                            naga::ImageClass::Storage {
+                                format: naga_format,
+                                access: naga_access,
+                            },
+                            usage,
+                        )
                     }
                     _ => return Err(BindingError::WrongType),
                 };
@@ -474,7 +491,7 @@ impl Resource {
             ResourceType::Buffer { size } => BindingType::Buffer {
                 ty: match self.class {
                     naga::StorageClass::Uniform => wgt::BufferBindingType::Uniform,
-                    naga::StorageClass::Storage => wgt::BufferBindingType::Storage {
+                    naga::StorageClass::Storage { .. } => wgt::BufferBindingType::Storage {
                         read_only: !shader_usage.contains(GlobalUse::WRITE),
                     },
                     _ => return Err(BindingError::WrongType),
@@ -517,7 +534,7 @@ impl Resource {
                         view_dimension,
                         multisampled: multi,
                     },
-                    naga::ImageClass::Storage(format) => BindingType::StorageTexture {
+                    naga::ImageClass::Storage { format, .. } => BindingType::StorageTexture {
                         access: if shader_usage == GlobalUse::WRITE || shader_usage.is_empty() {
                             wgt::StorageTextureAccess::WriteOnly
                         } else if !features

wgpu-hal/Cargo.toml

@@ -65,11 +65,11 @@ core-graphics-types = "0.1"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a7ac13a"
+rev = "300039e"
 
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a7ac13a"
+rev = "300039e"
 features = ["wgsl-in"]
 
 [dev-dependencies]

wgpu-hal/src/gles/device.rs

@@ -28,7 +28,7 @@ impl CompilationContext<'_> {
             }
             let register = match var.class {
                 naga::StorageClass::Uniform => super::BindingRegister::UniformBuffers,
-                naga::StorageClass::Storage => super::BindingRegister::StorageBuffers,
+                naga::StorageClass::Storage { .. } => super::BindingRegister::StorageBuffers,
                 _ => continue,
             };
 
@@ -48,10 +48,12 @@ impl CompilationContext<'_> {
 
         for (name, mapping) in reflection_info.texture_mapping {
             let var = &module.global_variables[mapping.texture];
-            let register = if var.storage_access.is_empty() {
-                super::BindingRegister::Textures
-            } else {
-                super::BindingRegister::Images
+            let register = match module.types[var.ty].inner {
+                naga::TypeInner::Image {
+                    class: naga::ImageClass::Storage { .. },
+                    ..
+                } => super::BindingRegister::Images,
+                _ => super::BindingRegister::Textures,
             };
 
             let tex_br = var.binding.as_ref().unwrap();

wgpu-hal/src/metal/device.rs

@@ -114,10 +114,14 @@ impl super::Device {
                     Some(ref br) => br.clone(),
                     None => continue,
                 };
-                // check for an immutable buffer
-                if !ep_info[var_handle].is_empty()
-                    && !var.storage_access.contains(naga::StorageAccess::STORE)
+                let storage_access_store = if let naga::StorageClass::Storage { access } = var.class
                 {
+                    access.contains(naga::StorageAccess::STORE)
+                } else {
+                    false
+                };
+                // check for an immutable buffer
+                if !ep_info[var_handle].is_empty() && !storage_access_store {
                     let psm = &layout.naga_options.per_stage_map[naga_stage];
                     let slot = psm.resources[&br].buffer.unwrap();
                     immutable_buffer_mask |= 1 << slot;

wgpu/Cargo.toml

@@ -75,13 +75,13 @@ env_logger = "0.8"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a7ac13a"
+rev = "300039e"
 optional = true
 
 # used to test all the example shaders
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a7ac13a"
+rev = "300039e"
 features = ["wgsl-in"]
 
 [[example]]

wgpu/examples/boids/compute.wgsl

@@ -20,8 +20,8 @@ struct Particles {
 };
 
 [[group(0), binding(0)]] var<uniform> params : SimParams;
-[[group(0), binding(1)]] var<storage> particlesSrc : [[access(read)]] Particles;
-[[group(0), binding(2)]] var<storage> particlesDst : [[access(read_write)]] Particles;
+[[group(0), binding(1)]] var<storage, read> particlesSrc : Particles;
+[[group(0), binding(2)]] var<storage, read_write> particlesDst : Particles;
 
 // https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
 [[stage(compute), workgroup_size(64)]]

wgpu/examples/hello-compute/shader.wgsl

@@ -4,7 +4,7 @@ struct PrimeIndices {
 }; // this is used as both input and output for convenience
 
 [[group(0), binding(0)]]
-var<storage> v_indices: [[access(read_write)]] PrimeIndices;
+var<storage, read_write> v_indices: PrimeIndices;
 
 // The Collatz Conjecture states that for any integer n:
 // If n is even, n = n/2

wgpu/examples/shadow/shader.wgsl

@@ -55,7 +55,7 @@ struct Lights {
 };
 
 [[group(0), binding(1)]]
-var<storage> s_lights: [[access(read)]] Lights;
+var<storage, read> s_lights: Lights;
 [[group(0), binding(2)]]
 var t_shadow: texture_depth_2d_array;
 [[group(0), binding(3)]]

wgpu/tests/vertex_indices/draw.vert.wgsl

@@ -4,7 +4,7 @@ struct Indices {
 }; // this is used as both input and output for convenience
 
 [[group(0), binding(0)]]
-var<storage> indices: [[access(read_write)]] Indices;
+var<storage, read_write> indices: Indices;
 
 [[stage(vertex)]]
 fn vs_main([[builtin(instance_index)]] instance: u32, [[builtin(vertex_index)]] index: u32) -> [[builtin(position)]] vec4<f32> {