[wgsl] Storage buffer/texture access (#1142)

* Resurrect texture_storage_* tests
* Test parsing of `var<storage,write>`
* Default storage textures to READ
* Restore default features
* Fix glsl/hlsl/msl/spv front and back ends
* Add missing test outputs
* All-around fixes for the storage access

Co-authored-by: Dzmitry Malyshau <kvarkus@gmail.com>

src/back/glsl/features.rs

@@ -298,7 +298,7 @@ impl<'a, W> Writer<'a, W> {
                                 self.features.request(Features::MULTISAMPLED_TEXTURE_ARRAYS);
                             }
                         }
-                        ImageClass::Storage(format) => match format {
+                        ImageClass::Storage { format, .. } => match format {
                             StorageFormat::R8Unorm
                             | StorageFormat::R8Snorm
                             | StorageFormat::R8Uint
@@ -336,7 +336,7 @@ impl<'a, W> Writer<'a, W> {
             }
             match global.class {
                 StorageClass::WorkGroup => self.features.request(Features::COMPUTE_SHADER),
-                StorageClass::Storage => self.features.request(Features::BUFFER_STORAGE),
+                StorageClass::Storage { .. } => self.features.request(Features::BUFFER_STORAGE),
                 StorageClass::PushConstant => return Err(Error::PushConstantNotSupported),
                 _ => {}
             }

src/back/glsl/mod.rs

@@ -497,11 +497,11 @@ impl<'a, W: Write> Writer<'a, W> {
                     class,
                 } => {
                     // Gather the storage format if needed
-                    let layout_storage_format = match self.module.types[global.ty].inner {
+                    let storage_format_access = match self.module.types[global.ty].inner {
                         TypeInner::Image {
-                            class: crate::ImageClass::Storage(format),
+                            class: crate::ImageClass::Storage { format, access },
                             ..
-                        } => Some(glsl_storage_format(format)),
+                        } => Some((format, access)),
                         _ => None,
                     };
                     // Gether the location if needed
@@ -513,23 +513,24 @@ impl<'a, W: Write> Writer<'a, W> {
                     };
 
                     // Write all the layout qualifiers
-                    if layout_binding.is_some() || layout_storage_format.is_some() {
+                    if layout_binding.is_some() || storage_format_access.is_some() {
                         write!(self.out, "layout(")?;
                         if let Some(binding) = layout_binding {
                             write!(self.out, "binding = {}", binding)?;
                         }
-                        if let Some(format) = layout_storage_format {
+                        if let Some((format, _)) = storage_format_access {
+                            let format_str = glsl_storage_format(format);
                             let separator = match layout_binding {
                                 Some(_) => ",",
                                 None => "",
                             };
-                            write!(self.out, "{}{}", separator, format)?;
+                            write!(self.out, "{}{}", separator, format_str)?;
                         }
                         write!(self.out, ") ")?;
                     }
 
-                    if let Some(storage_access) = glsl_storage_access(global.storage_access) {
-                        write!(self.out, "{} ", storage_access)?;
+                    if let Some((_, access)) = storage_format_access {
+                        self.write_storage_access(access)?;
                     }
 
                     // All images in glsl are `uniform`
@@ -760,7 +761,7 @@ impl<'a, W: Write> Writer<'a, W> {
             Ic::Sampled { kind, multi: false } => ("sampler", kind, "", ""),
             Ic::Depth { multi: true } => ("sampler", crate::ScalarKind::Float, "MS", ""),
             Ic::Depth { multi: false } => ("sampler", crate::ScalarKind::Float, "", "Shadow"),
-            Ic::Storage(format) => ("image", format.into(), "", ""),
+            Ic::Storage { format, .. } => ("image", format.into(), "", ""),
         };
 
         write!(
@@ -798,8 +799,8 @@ impl<'a, W: Write> Writer<'a, W> {
             }
         }
 
-        if let Some(storage_access) = glsl_storage_access(global.storage_access) {
-            write!(self.out, "{} ", storage_access)?;
+        if let crate::StorageClass::Storage { access } = global.class {
+            self.write_storage_access(access)?;
         }
 
         // Write the storage class
@@ -1013,7 +1014,7 @@ impl<'a, W: Write> Writer<'a, W> {
                 } => {
                     // Write the storage format if needed
                     if let TypeInner::Image {
-                        class: crate::ImageClass::Storage(format),
+                        class: crate::ImageClass::Storage { format, .. },
                         ..
                     } = this.module.types[arg.ty].inner
                     {
@@ -1939,7 +1940,7 @@ impl<'a, W: Write> Writer<'a, W> {
 
                 let fun_name = match class {
                     crate::ImageClass::Sampled { .. } => "texelFetch",
-                    crate::ImageClass::Storage(_) => "imageLoad",
+                    crate::ImageClass::Storage { .. } => "imageLoad",
                     // TODO: Is there even a function for this?
                     crate::ImageClass::Depth { multi: _ } => {
                         return Err(Error::Custom("TODO: depth sample loads".to_string()))
@@ -1992,7 +1993,7 @@ impl<'a, W: Write> Writer<'a, W> {
                                     write!(self.out, "0",)?;
                                 }
                             }
-                            ImageClass::Storage(_) => {
+                            ImageClass::Storage { .. } => {
                                 write!(self.out, "imageSize(")?;
                                 self.write_expr(image, ctx)?;
                             }
@@ -2007,7 +2008,7 @@ impl<'a, W: Write> Writer<'a, W> {
                     crate::ImageQuery::NumLayers => {
                         let fun_name = match class {
                             ImageClass::Sampled { .. } | ImageClass::Depth { .. } => "textureSize",
-                            ImageClass::Storage(_) => "imageSize",
+                            ImageClass::Storage { .. } => "imageSize",
                         };
                         write!(self.out, "{}(", fun_name)?;
                         self.write_expr(image, ctx)?;
@@ -2019,7 +2020,7 @@ impl<'a, W: Write> Writer<'a, W> {
                             ImageClass::Sampled { .. } | ImageClass::Depth { .. } => {
                                 "textureSamples"
                             }
-                            ImageClass::Storage(_) => "imageSamples",
+                            ImageClass::Storage { .. } => "imageSamples",
                         };
                         write!(self.out, "{}(", fun_name)?;
                         self.write_expr(image, ctx)?;
@@ -2417,6 +2418,21 @@ impl<'a, W: Write> Writer<'a, W> {
         Ok(())
     }
 
+    /// Helper function that return the glsl storage access string of [`StorageAccess`](crate::StorageAccess)
+    ///
+    /// glsl allows adding both `readonly` and `writeonly` but this means that
+    /// they can only be used to query information about the resource which isn't what
+    /// we want here so when storage access is both `LOAD` and `STORE` add no modifiers
+    fn write_storage_access(&mut self, storage_access: crate::StorageAccess) -> BackendResult {
+        if !storage_access.contains(crate::StorageAccess::STORE) {
+            write!(self.out, "readonly ")?;
+        }
+        if !storage_access.contains(crate::StorageAccess::LOAD) {
+            write!(self.out, "writeonly ")?;
+        }
+        Ok(())
+    }
+
     /// Helper method used to produce the reflection info that's returned to the user
     ///
     /// It takes an iterator of [`Function`](crate::Function) references instead of
@@ -2454,7 +2470,7 @@ impl<'a, W: Write> Writer<'a, W> {
             }
             match self.module.types[var.ty].inner {
                 crate::TypeInner::Struct { .. } => match var.class {
-                    crate::StorageClass::Uniform | crate::StorageClass::Storage => {
+                    crate::StorageClass::Uniform | crate::StorageClass::Storage { .. } => {
                         let name = self.reflection_names[&var.ty].clone();
                         uniforms.insert(handle, name);
                     }
@@ -2568,7 +2584,7 @@ fn glsl_storage_class(class: crate::StorageClass) -> Option<&'static str> {
     match class {
         Sc::Function => None,
         Sc::Private => None,
-        Sc::Storage => Some("buffer"),
+        Sc::Storage { .. } => Some("buffer"),
         Sc::Uniform => Some("uniform"),
         Sc::Handle => Some("uniform"),
         Sc::WorkGroup => Some("shared"),
@@ -2650,21 +2666,6 @@ fn glsl_storage_format(format: crate::StorageFormat) -> &'static str {
     }
 }
 
-/// Helper function that return the glsl storage access string of [`StorageAccess`](crate::StorageAccess)
-///
-/// glsl allows adding both `readonly` and `writeonly` but this means that
-/// they can only be used to query information about the resource which isn't what
-/// we want here so when storage access is both `LOAD` and `STORE` add no modifiers
-fn glsl_storage_access(storage_access: crate::StorageAccess) -> Option<&'static str> {
-    if storage_access == crate::StorageAccess::LOAD {
-        Some("readonly")
-    } else if storage_access == crate::StorageAccess::STORE {
-        Some("writeonly")
-    } else {
-        None
-    }
-}
-
 fn is_value_init_supported(module: &crate::Module, ty: Handle<crate::Type>) -> bool {
     match module.types[ty].inner {
         TypeInner::Scalar { .. } | TypeInner::Vector { .. } | TypeInner::Matrix { .. } => true,

src/back/hlsl/help.rs

@@ -105,9 +105,13 @@ impl<'a, W: Write> super::Writer<'a, W> {
         arrayed: bool,
         class: crate::ImageClass,
     ) -> BackendResult {
+        let access_str = match class {
+            crate::ImageClass::Storage { .. } => "RW",
+            _ => "",
+        };
         let dim_str = dim.to_hlsl_str();
         let arrayed_str = if arrayed { "Array" } else { "" };
-        write!(self.out, "Texture{}{}", dim_str, arrayed_str)?;
+        write!(self.out, "{}Texture{}{}", access_str, dim_str, arrayed_str)?;
         match class {
             crate::ImageClass::Depth { multi } => {
                 let multi_str = if multi { "MS" } else { "" };
@@ -118,7 +122,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
                 let scalar_kind_str = kind.to_hlsl_str(4)?;
                 write!(self.out, "{}<{}4>", multi_str, scalar_kind_str)?
             }
-            crate::ImageClass::Storage(format) => {
+            crate::ImageClass::Storage { format, .. } => {
                 let storage_format_str = format.to_hlsl_str();
                 write!(self.out, "<{}>", storage_format_str)?
             }
@@ -195,9 +199,8 @@ impl<'a, W: Write> super::Writer<'a, W> {
             crate::ImageClass::Sampled { multi: true, .. } => "MS",
             crate::ImageClass::Depth { multi: true } => "DepthMS",
             crate::ImageClass::Depth { multi: false } => "Depth",
-            crate::ImageClass::Sampled { multi: false, .. } | crate::ImageClass::Storage { .. } => {
-                ""
-            }
+            crate::ImageClass::Sampled { multi: false, .. } => "",
+            crate::ImageClass::Storage { .. } => "RW",
         };
         let arrayed_str = if query.arrayed { "Array" } else { "" };
         let query_str = match query.query {
@@ -226,7 +229,10 @@ impl<'a, W: Write> super::Writer<'a, W> {
         expr_handle: Handle<crate::Expression>,
         func_ctx: &FunctionCtx,
     ) -> BackendResult {
-        use crate::{back::INDENT, ImageDimension as IDim};
+        use crate::{
+            back::{COMPONENTS, INDENT},
+            ImageDimension as IDim,
+        };
 
         const ARGUMENT_VARIABLE_NAME: &str = "texture";
         const RETURN_VARIABLE_NAME: &str = "ret";
@@ -253,15 +259,24 @@ impl<'a, W: Write> super::Writer<'a, W> {
         writeln!(self.out, "{{")?;
 
         let array_coords = if wiq.arrayed { 1 } else { 0 };
+        // extra parameter is the mip level count or the sample count
+        let extra_coords = match wiq.class {
+            crate::ImageClass::Storage { .. } => 0,
+            crate::ImageClass::Sampled { .. } | crate::ImageClass::Depth { .. } => 1,
+        };
+
         // GetDimensions Overloaded Methods
         // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions#overloaded-methods
         let (ret_swizzle, number_of_params) = match wiq.query {
-            ImageQuery::Size | ImageQuery::SizeLevel => match wiq.dim {
-                IDim::D1 => ("x", 1 + array_coords),
-                IDim::D2 => ("xy", 3 + array_coords),
-                IDim::D3 => ("xyz", 4),
-                IDim::Cube => ("xy", 3 + array_coords),
-            },
+            ImageQuery::Size | ImageQuery::SizeLevel => {
+                let ret = match wiq.dim {
+                    IDim::D1 => "x",
+                    IDim::D2 => "xy",
+                    IDim::D3 => "xyz",
+                    IDim::Cube => "xy",
+                };
+                (ret, ret.len() + array_coords + extra_coords)
+            }
             ImageQuery::NumLevels | ImageQuery::NumSamples | ImageQuery::NumLayers => {
                 if wiq.arrayed || wiq.dim == IDim::D3 {
                     ("w", 4)
@@ -284,18 +299,16 @@ impl<'a, W: Write> super::Writer<'a, W> {
             }
             _ => match wiq.class {
                 crate::ImageClass::Sampled { multi: true, .. }
-                | crate::ImageClass::Depth { multi: true } => {}
-                _ => match wiq.dim {
+                | crate::ImageClass::Depth { multi: true }
+                | crate::ImageClass::Storage { .. } => {}
+                _ => {
                     // Write zero mipmap level for supported types
-                    IDim::D2 | IDim::D3 | IDim::Cube => {
-                        write!(self.out, "0, ")?;
-                    }
-                    IDim::D1 => {}
-                },
+                    write!(self.out, "0, ")?;
+                }
             },
         }
 
-        for component in crate::back::COMPONENTS[..number_of_params - 1].iter() {
+        for component in COMPONENTS[..number_of_params - 1].iter() {
             write!(self.out, "{}.{}, ", RETURN_VARIABLE_NAME, component)?;
         }
 
@@ -304,7 +317,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
             self.out,
             "{}.{}",
             RETURN_VARIABLE_NAME,
-            crate::back::COMPONENTS[number_of_params - 1]
+            COMPONENTS[number_of_params - 1]
         )?;
 
         writeln!(self.out, ");")?;
@@ -344,10 +357,12 @@ impl<'a, W: Write> super::Writer<'a, W> {
                         }
                         ref other => unreachable!("Array length of base {:?}", other),
                     };
+                    let storage_access = match global_var.class {
+                        crate::StorageClass::Storage { access } => access,
+                        _ => crate::StorageAccess::default(),
+                    };
                     let wal = WrappedArrayLength {
-                        writable: global_var
-                            .storage_access
-                            .contains(crate::StorageAccess::STORE),
+                        writable: storage_access.contains(crate::StorageAccess::STORE),
                     };
 
                     if !self.wrapped_array_lengths.contains(&wal) {

src/back/hlsl/writer.rs

@@ -396,24 +396,24 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 write!(self.out, "cbuffer")?;
                 "b"
             }
-            crate::StorageClass::Storage => {
-                let (prefix, register) =
-                    if global.storage_access.contains(crate::StorageAccess::STORE) {
-                        ("RW", "u")
-                    } else {
-                        ("", "t")
-                    };
+            crate::StorageClass::Storage { access } => {
+                let (prefix, register) = if access.contains(crate::StorageAccess::STORE) {
+                    ("RW", "u")
+                } else {
+                    ("", "t")
+                };
                 write!(self.out, "{}ByteAddressBuffer", prefix)?;
                 register
             }
             crate::StorageClass::Handle => {
-                let register = if let TypeInner::Sampler { .. } = *inner {
-                    "s"
-                } else if global.storage_access.contains(crate::StorageAccess::STORE) {
-                    write!(self.out, "RW")?;
-                    "u"
-                } else {
-                    "t"
+                let register = match *inner {
+                    TypeInner::Sampler { .. } => "s",
+                    // all storage textures are UAV, unconditionally
+                    TypeInner::Image {
+                        class: crate::ImageClass::Storage { .. },
+                        ..
+                    } => "u",
+                    _ => "t",
                 };
                 self.write_type(module, global.ty)?;
                 register
@@ -954,11 +954,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     _ => None,
                 };
 
-                if func_ctx.info[pointer]
+                if let Some(crate::StorageClass::Storage { .. }) = func_ctx.info[pointer]
                     .ty
                     .inner_with(&module.types)
                     .pointer_class()
-                    == Some(crate::StorageClass::Storage)
                 {
                     let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
                     self.write_storage_store(
@@ -1208,11 +1207,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 write!(self.out, ")")?;
             }
             Expression::Access { base, index } => {
-                if func_ctx.info[expr]
+                if let Some(crate::StorageClass::Storage { .. }) = func_ctx.info[expr]
                     .ty
                     .inner_with(&module.types)
                     .pointer_class()
-                    == Some(crate::StorageClass::Storage)
                 {
                     // do nothing, the chain is written on `Load`/`Store`
                 } else {
@@ -1223,11 +1221,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 }
             }
             Expression::AccessIndex { base, index } => {
-                if func_ctx.info[expr]
+                if let Some(crate::StorageClass::Storage { .. }) = func_ctx.info[expr]
                     .ty
                     .inner_with(&module.types)
                     .pointer_class()
-                    == Some(crate::StorageClass::Storage)
                 {
                     // do nothing, the chain is written on `Load`/`Store`
                 } else {
@@ -1389,22 +1386,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 index,
             } => {
                 // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-load
-                let ms = match *func_ctx.info[image].ty.inner_with(&module.types) {
-                    TypeInner::Image {
-                        class: crate::ImageClass::Sampled { multi, .. },
-                        ..
-                    }
-                    | TypeInner::Image {
-                        class: crate::ImageClass::Depth { multi },
-                        ..
-                    } => multi,
-                    _ => false,
+                let (ms, storage) = match *func_ctx.info[image].ty.inner_with(&module.types) {
+                    TypeInner::Image { class, .. } => match class {
+                        crate::ImageClass::Sampled { multi, .. }
+                        | crate::ImageClass::Depth { multi } => (multi, false),
+                        crate::ImageClass::Storage { .. } => (false, true),
+                    },
+                    _ => (false, false),
                 };
 
                 self.write_expr(module, image, func_ctx)?;
                 write!(self.out, ".Load(")?;
 
-                let mip_level = if ms {
+                let mip_level = if ms || storage {
                     MipLevelCoordinate::NotApplicable
                 } else {
                     match index {
@@ -1436,27 +1430,30 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     write!(self.out, ".x")?;
                 }
             }
-            Expression::GlobalVariable(handle) => {
-                if module.global_variables[handle].class != crate::StorageClass::Storage {
+            Expression::GlobalVariable(handle) => match module.global_variables[handle].class {
+                crate::StorageClass::Storage { .. } => {}
+                _ => {
                     let name = &self.names[&NameKey::GlobalVariable(handle)];
                     write!(self.out, "{}", name)?;
                 }
-            }
+            },
             Expression::LocalVariable(handle) => {
                 write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
             }
             Expression::Load { pointer } => {
-                if func_ctx.info[pointer]
+                match func_ctx.info[pointer]
                     .ty
                     .inner_with(&module.types)
                     .pointer_class()
-                    == Some(crate::StorageClass::Storage)
                 {
-                    let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
-                    let result_ty = func_ctx.info[expr].ty.clone();
-                    self.write_storage_load(module, var_handle, result_ty, func_ctx)?;
-                } else {
-                    self.write_expr(module, pointer, func_ctx)?;
+                    Some(crate::StorageClass::Storage { .. }) => {
+                        let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
+                        let result_ty = func_ctx.info[expr].ty.clone();
+                        self.write_storage_load(module, var_handle, result_ty, func_ctx)?;
+                    }
+                    _ => {
+                        self.write_expr(module, pointer, func_ctx)?;
+                    }
                 }
             }
             Expression::Unary { op, expr } => {
@@ -1613,8 +1610,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     _ => unreachable!(),
                 };
 
+                let storage_access = match var.class {
+                    crate::StorageClass::Storage { access } => access,
+                    _ => crate::StorageAccess::default(),
+                };
                 let wrapped_array_length = WrappedArrayLength {
-                    writable: var.storage_access.contains(crate::StorageAccess::STORE),
+                    writable: storage_access.contains(crate::StorageAccess::STORE),
                 };
 
                 write!(self.out, "((")?;

src/back/msl/writer.rs

@@ -81,7 +81,7 @@ impl<'a> Display for TypeContext<'a> {
                     first_time: false,
                     ..*self
                 };
-                let class_name = match class.get_name(self.access) {
+                let class_name = match class.to_msl_name() {
                     Some(name) => name,
                     None => return Ok(()),
                 };
@@ -93,7 +93,7 @@ impl<'a> Display for TypeContext<'a> {
                 width: _,
                 class,
             } => {
-                let class_name = match class.get_name(self.access) {
+                let class_name = match class.to_msl_name() {
                     Some(name) => name,
                     None => return Ok(()),
                 };
@@ -105,7 +105,7 @@ impl<'a> Display for TypeContext<'a> {
                 width: _,
                 class,
             } => {
-                let class_name = match class.get_name(self.access) {
+                let class_name = match class.to_msl_name() {
                     Some(name) => name,
                     None => return Ok(()),
                 };
@@ -157,7 +157,7 @@ impl<'a> Display for TypeContext<'a> {
                         };
                         ("depth", msaa_str, crate::ScalarKind::Float, access)
                     }
-                    crate::ImageClass::Storage(format) => {
+                    crate::ImageClass::Storage { format, .. } => {
                         let access = if self
                             .access
                             .contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE)
@@ -207,15 +207,26 @@ impl<'a> TypedGlobalVariable<'a> {
     fn try_fmt<W: Write>(&self, out: &mut W) -> BackendResult {
         let var = &self.module.global_variables[self.handle];
         let name = &self.names[&NameKey::GlobalVariable(self.handle)];
+
+        let storage_access = match var.class {
+            crate::StorageClass::Storage { access } => access,
+            _ => match self.module.types[var.ty].inner {
+                crate::TypeInner::Image {
+                    class: crate::ImageClass::Storage { access, .. },
+                    ..
+                } => access,
+                _ => crate::StorageAccess::default(),
+            },
+        };
         let ty_name = TypeContext {
             handle: var.ty,
             arena: &self.module.types,
             names: self.names,
-            access: var.storage_access,
+            access: storage_access,
             first_time: false,
         };
 
-        let (space, access, reference) = match var.class.get_name(var.storage_access) {
+        let (space, access, reference) = match var.class.to_msl_name() {
             Some(space) if self.reference => {
                 let access = match var.class {
                     crate::StorageClass::Private | crate::StorageClass::WorkGroup
@@ -366,7 +377,7 @@ impl crate::StorageClass {
     fn needs_pass_through(&self) -> bool {
         match *self {
             crate::StorageClass::Uniform
-            | crate::StorageClass::Storage
+            | crate::StorageClass::Storage { .. }
             | crate::StorageClass::Private
             | crate::StorageClass::PushConstant
             | crate::StorageClass::Handle => true,
@@ -374,16 +385,14 @@ impl crate::StorageClass {
         }
     }
 
-    fn get_name(&self, access: crate::StorageAccess) -> Option<&'static str> {
-        match *self {
+    fn to_msl_name(self) -> Option<&'static str> {
+        match self {
             Self::Handle => None,
             Self::Uniform | Self::PushConstant => Some("constant"),
-            //TODO: should still be "constant" for read-only buffers
-            Self::Storage => Some(if access.contains(crate::StorageAccess::STORE) {
-                "device"
-            } else {
-                "constant"
-            }),
+            Self::Storage { access } if access.contains(crate::StorageAccess::STORE) => {
+                Some("device")
+            }
+            Self::Storage { .. } => Some("constant"),
             Self::Private | Self::Function => Some("thread"),
             Self::WorkGroup => Some("threadgroup"),
         }

src/back/spv/block.rs

@@ -1043,7 +1043,7 @@ impl<'w> BlockContext<'w> {
                         };
 
                         let (query_op, level_id) = match class {
-                            Ic::Storage(_) => (spirv::Op::ImageQuerySize, None),
+                            Ic::Storage { .. } => (spirv::Op::ImageQuerySize, None),
                             _ => {
                                 let level_id = match level {
                                     Some(expr) => self.cached[expr],

src/back/spv/helpers.rs

@@ -25,7 +25,7 @@ pub(super) fn map_storage_class(class: crate::StorageClass) -> spirv::StorageCla
         crate::StorageClass::Handle => spirv::StorageClass::UniformConstant,
         crate::StorageClass::Function => spirv::StorageClass::Function,
         crate::StorageClass::Private => spirv::StorageClass::Private,
-        crate::StorageClass::Storage => spirv::StorageClass::StorageBuffer,
+        crate::StorageClass::Storage { .. } => spirv::StorageClass::StorageBuffer,
         crate::StorageClass::Uniform => spirv::StorageClass::Uniform,
         crate::StorageClass::WorkGroup => spirv::StorageClass::Workgroup,
         crate::StorageClass::PushConstant => spirv::StorageClass::PushConstant,

src/back/spv/instructions.rs

@@ -234,7 +234,7 @@ impl super::Instruction {
         let (depth, multi, sampled) = match image_class {
             crate::ImageClass::Sampled { kind: _, multi } => (false, multi, true),
             crate::ImageClass::Depth { multi } => (true, multi, true),
-            crate::ImageClass::Storage(_) => (false, false, false),
+            crate::ImageClass::Storage { .. } => (false, false, false),
         };
         instruction.add_operand(depth as u32);
         instruction.add_operand(arrayed as u32);
@@ -242,7 +242,7 @@ impl super::Instruction {
         instruction.add_operand(if sampled { 1 } else { 2 });
 
         let format = match image_class {
-            crate::ImageClass::Storage(format) => match format {
+            crate::ImageClass::Storage { format, .. } => match format {
                 crate::StorageFormat::R8Unorm => spirv::ImageFormat::R8,
                 crate::StorageFormat::R8Snorm => spirv::ImageFormat::R8Snorm,
                 crate::StorageFormat::R8Uint => spirv::ImageFormat::R8ui,

src/back/spv/writer.rs

@@ -702,7 +702,7 @@ impl Writer {
                 let kind = match class {
                     crate::ImageClass::Sampled { kind, multi: _ } => kind,
                     crate::ImageClass::Depth { multi: _ } => crate::ScalarKind::Float,
-                    crate::ImageClass::Storage(format) => {
+                    crate::ImageClass::Storage { format, .. } => {
                         let required_caps: &[_] = match dim {
                             crate::ImageDimension::D1 => &[spirv::Capability::Image1D],
                             crate::ImageDimension::Cube => &[spirv::Capability::ImageCubeArray],
@@ -1057,13 +1057,23 @@ impl Writer {
 
         use spirv::Decoration;
 
-        let access_decoration = match global_variable.storage_access {
-            crate::StorageAccess::LOAD => Some(Decoration::NonWritable),
-            crate::StorageAccess::STORE => Some(Decoration::NonReadable),
-            _ => None,
+        let storage_access = match global_variable.class {
+            crate::StorageClass::Storage { access } => Some(access),
+            _ => match ir_module.types[global_variable.ty].inner {
+                crate::TypeInner::Image {
+                    class: crate::ImageClass::Storage { access, .. },
+                    ..
+                } => Some(access),
+                _ => None,
+            },
         };
-        if let Some(decoration) = access_decoration {
-            self.decorate(id, decoration, &[]);
+        if let Some(storage_access) = storage_access {
+            if !storage_access.contains(crate::StorageAccess::LOAD) {
+                self.decorate(id, Decoration::NonReadable, &[]);
+            }
+            if !storage_access.contains(crate::StorageAccess::STORE) {
+                self.decorate(id, Decoration::NonWritable, &[]);
+            }
         }
 
         if let Some(ref res_binding) = global_variable.binding {
@@ -1104,10 +1114,14 @@ impl Writer {
         ir_module: &crate::Module,
         mod_info: &ModuleInfo,
     ) -> Result<(), Error> {
-        let has_storage_buffers = ir_module
-            .global_variables
-            .iter()
-            .any(|(_, var)| var.class == crate::StorageClass::Storage);
+        let has_storage_buffers =
+            ir_module
+                .global_variables
+                .iter()
+                .any(|(_, var)| match var.class {
+                    crate::StorageClass::Storage { .. } => true,
+                    _ => false,
+                });
         if self.physical_layout.version < 0x10300 && has_storage_buffers {
             // enable the storage buffer class on < SPV-1.3
             Instruction::extension("SPV_KHR_storage_buffer_storage_class")

src/back/wgsl/writer.rs

@@ -12,7 +12,6 @@ type BackendResult = Result<(), Error>;
 /// WGSL attribute
 /// https://gpuweb.github.io/gpuweb/wgsl/#attributes
 enum Attribute {
-    Access(crate::StorageAccess),
     Binding(u32),
     Block,
     BuiltIn(crate::BuiltIn),
@@ -297,16 +296,6 @@ impl<W: Write> Writer<W> {
         let mut attributes_str = String::new();
         for (index, attribute) in attributes.iter().enumerate() {
             let attribute_str = match *attribute {
-                Attribute::Access(access) => {
-                    let access_str = if access.is_all() {
-                        "read_write"
-                    } else if access.contains(crate::StorageAccess::LOAD) {
-                        "read"
-                    } else {
-                        "write"
-                    };
-                    format!("access({})", access_str)
-                }
                 Attribute::Block => String::from("block"),
                 Attribute::Location(id) => format!("location({})", id),
                 Attribute::BuiltIn(builtin_attrib) => {
@@ -491,10 +480,14 @@ impl<W: Write> Writer<W> {
                         if multi { "multisampled_" } else { "" },
                         String::from(""),
                     ),
-                    Ic::Storage(storage_format) => (
+                    Ic::Storage { format, access } => (
                         "storage_",
                         "",
-                        format!("<{}>", storage_format_str(storage_format)),
+                        if access.contains(crate::StorageAccess::STORE) {
+                            format!("<{},write>", storage_format_str(format))
+                        } else {
+                            format!("<{}>", storage_format_str(format))
+                        },
                     ),
                 };
                 let ty_str = format!(
@@ -1329,10 +1322,7 @@ impl<W: Write> Writer<W> {
             write!(self.out, "<{}>", storage_class)?;
         }
         write!(self.out, " {}: ", name)?;
-        // Write access attribute if present
-        if !global.storage_access.is_empty() {
-            self.write_attributes(&[Attribute::Access(global.storage_access)], true)?;
-        }
+
         // Write global type
         self.write_type(module, global.ty)?;
 
@@ -1586,7 +1576,11 @@ fn storage_class_str(storage_class: crate::StorageClass) -> Option<&'static str>
     match storage_class {
         Sc::Private => Some("private"),
         Sc::Uniform => Some("uniform"),
-        Sc::Storage => Some("storage"),
+        Sc::Storage { access } => Some(if access.contains(crate::StorageAccess::STORE) {
+            "storage,read_write"
+        } else {
+            "storage"
+        }),
         Sc::PushConstant => Some("push_constant"),
         Sc::WorkGroup => Some("workgroup"),
         Sc::Function | Sc::Handle => None,

src/front/glsl/parser.rs

@@ -241,7 +241,9 @@ impl<'source, 'program, 'options> Parser<'source, 'program, 'options> {
                         StorageQualifier::StorageClass(StorageClass::Uniform),
                     ),
                     TokenValue::Buffer => TypeQualifier::StorageQualifier(
-                        StorageQualifier::StorageClass(StorageClass::Storage),
+                        StorageQualifier::StorageClass(StorageClass::Storage {
+                            access: crate::StorageAccess::default(),
+                        }),
                     ),
                     TokenValue::Sampling(s) => TypeQualifier::Sampling(s),
                     TokenValue::PrecisionQualifier(p) => TypeQualifier::Precision(p),

src/front/glsl/variables.rs

@@ -1,7 +1,7 @@
 use crate::{
-    Binding, Block, BuiltIn, Constant, Expression, GlobalVariable, Handle, ImageClass,
-    Interpolation, LocalVariable, ScalarKind, StorageAccess, StorageClass, SwizzleComponent, Type,
-    TypeInner, VectorSize,
+    Binding, Block, BuiltIn, Constant, Expression, GlobalVariable, Handle, Interpolation,
+    LocalVariable, ScalarKind, StorageAccess, StorageClass, SwizzleComponent, Type, TypeInner,
+    VectorSize,
 };
 
 use super::ast::*;
@@ -57,7 +57,6 @@ impl Program<'_> {
                 binding: None,
                 ty,
                 init: None,
-                storage_access: StorageAccess::empty(),
             });
 
             let idx = self.entry_args.len();
@@ -407,7 +406,7 @@ impl Program<'_> {
             match storage {
                 StorageQualifier::StorageClass(StorageClass::PushConstant)
                 | StorageQualifier::StorageClass(StorageClass::Uniform)
-                | StorageQualifier::StorageClass(StorageClass::Storage) => {}
+                | StorageQualifier::StorageClass(StorageClass::Storage { .. }) => {}
                 _ => {
                     return Err(ErrorKind::SemanticError(
                         meta,
@@ -445,7 +444,6 @@ impl Program<'_> {
                 binding: None,
                 ty,
                 init,
-                storage_access: StorageAccess::empty(),
             });
 
             let idx = self.entry_args.len();
@@ -490,29 +488,20 @@ impl Program<'_> {
             return Ok(GlobalOrConstant::Constant(init));
         }
 
-        let (class, storage_access) = match self.module.types[ty].inner {
-            TypeInner::Image { class, .. } => (
-                StorageClass::Handle,
-                if let ImageClass::Storage(_) = class {
-                    // TODO: Add support for qualifiers such as readonly,
-                    // writeonly and readwrite
-                    StorageAccess::all()
-                } else {
-                    StorageAccess::empty()
-                },
-            ),
-            TypeInner::Sampler { .. } => (StorageClass::Handle, StorageAccess::empty()),
+        // TODO: Add support for qualifiers such as readonly, writeonly and readwrite
+        let class = match self.module.types[ty].inner {
+            TypeInner::Image { .. } => StorageClass::Handle,
+            TypeInner::Sampler { .. } => StorageClass::Handle,
             _ => {
-                if let StorageQualifier::StorageClass(StorageClass::Storage) = storage {
-                    (StorageClass::Storage, StorageAccess::all())
+                if let StorageQualifier::StorageClass(StorageClass::Storage { .. }) = storage {
+                    StorageClass::Storage {
+                        access: StorageAccess::all(),
+                    }
                 } else {
-                    (
-                        match storage {
-                            StorageQualifier::StorageClass(class) => class,
-                            _ => StorageClass::Private,
-                        },
-                        StorageAccess::empty(),
-                    )
+                    match storage {
+                        StorageQualifier::StorageClass(class) => class,
+                        _ => StorageClass::Private,
+                    }
                 }
             }
         };
@@ -523,7 +512,6 @@ impl Program<'_> {
             binding,
             ty,
             init,
-            storage_access,
         });
 
         if let Some(name) = name {

src/front/spv/convert.rs

@@ -154,7 +154,9 @@ pub(super) fn map_storage_class(word: spirv::Word) -> Result<super::ExtendedClas
         Some(Sc::Output) => Ec::Output,
         Some(Sc::Private) => Ec::Global(crate::StorageClass::Private),
         Some(Sc::UniformConstant) => Ec::Global(crate::StorageClass::Handle),
-        Some(Sc::StorageBuffer) => Ec::Global(crate::StorageClass::Storage),
+        Some(Sc::StorageBuffer) => Ec::Global(crate::StorageClass::Storage {
+            access: crate::StorageAccess::default(),
+        }),
         // we expect the `Storage` case to be filtered out before calling this function.
         Some(Sc::Uniform) => Ec::Global(crate::StorageClass::Uniform),
         Some(Sc::Workgroup) => Ec::Global(crate::StorageClass::WorkGroup),

src/front/spv/mod.rs

@@ -2912,7 +2912,9 @@ impl<I: Iterator<Item = u32>> Parser<I> {
             .lookup_storage_buffer_types
             .contains_key(&base_lookup_ty.handle)
         {
-            crate::StorageClass::Storage
+            crate::StorageClass::Storage {
+                access: crate::StorageAccess::default(),
+            }
         } else {
             match map_storage_class(storage_class)? {
                 ExtendedClass::Global(class) => class,
@@ -3161,7 +3163,10 @@ impl<I: Iterator<Item = u32>> Parser<I> {
 
         let inner = crate::TypeInner::Image {
             class: if format != 0 {
-                crate::ImageClass::Storage(map_image_format(format)?)
+                crate::ImageClass::Storage {
+                    format: map_image_format(format)?,
+                    access: crate::StorageAccess::default(),
+                }
             } else {
                 crate::ImageClass::Sampled {
                     kind,
@@ -3418,18 +3423,25 @@ impl<I: Iterator<Item = u32>> Parser<I> {
 
         let original_ty = self.lookup_type.lookup(type_id)?.handle;
         let (effective_ty, is_storage) = match module.types[original_ty].inner {
-            crate::TypeInner::Pointer { base, class } => {
-                (base, class == crate::StorageClass::Storage)
-            }
+            crate::TypeInner::Pointer { base, class } => (
+                base,
+                match class {
+                    crate::StorageClass::Storage { .. } => true,
+                    _ => false,
+                },
+            ),
             crate::TypeInner::Image {
-                class: crate::ImageClass::Storage(_),
+                class: crate::ImageClass::Storage { .. },
                 ..
             } => (original_ty, true),
             _ => (original_ty, false),
         };
         let (ext_class, type_storage_access) =
             match self.lookup_storage_buffer_types.get(&effective_ty) {
-                Some(&access) => (ExtendedClass::Global(crate::StorageClass::Storage), access),
+                Some(&access) => (
+                    ExtendedClass::Global(crate::StorageClass::Storage { access }),
+                    access,
+                ),
                 None => (
                     map_storage_class(storage_class)?,
                     crate::StorageAccess::all(),
@@ -3449,7 +3461,8 @@ impl<I: Iterator<Item = u32>> Parser<I> {
 
         let (inner, var) = match ext_class {
             ExtendedClass::Global(class) => {
-                let storage_access = if is_storage {
+                // TODO actually apply decorators
+                let _storage_access = if is_storage {
                     let mut access = type_storage_access;
                     if dec.flags.contains(DecorationFlags::NON_READABLE) {
                         access &= !crate::StorageAccess::LOAD;
@@ -3468,7 +3481,6 @@ impl<I: Iterator<Item = u32>> Parser<I> {
                     class,
                     ty: effective_ty,
                     init,
-                    storage_access,
                 };
                 (Variable::Global, var)
             }
@@ -3513,7 +3525,6 @@ impl<I: Iterator<Item = u32>> Parser<I> {
                     binding: None,
                     ty: effective_ty,
                     init: None,
-                    storage_access: crate::StorageAccess::empty(),
                 };
                 let inner = Variable::Input(crate::FunctionArgument {
                     name: dec.name,
@@ -3584,7 +3595,6 @@ impl<I: Iterator<Item = u32>> Parser<I> {
                     binding: None,
                     ty: effective_ty,
                     init,
-                    storage_access: crate::StorageAccess::empty(),
                 };
                 let inner = Variable::Output(crate::FunctionResult {
                     ty: effective_ty,

src/front/wgsl/conv.rs

@@ -5,7 +5,9 @@ pub fn map_storage_class(word: &str, span: Span) -> Result<crate::StorageClass,
         "private" => Ok(crate::StorageClass::Private),
         "workgroup" => Ok(crate::StorageClass::WorkGroup),
         "uniform" => Ok(crate::StorageClass::Uniform),
-        "storage" => Ok(crate::StorageClass::Storage),
+        "storage" => Ok(crate::StorageClass::Storage {
+            access: crate::StorageAccess::default(),
+        }),
         "push_constant" => Ok(crate::StorageClass::PushConstant),
         _ => Err(Error::UnknownStorageClass(span)),
     }

src/front/wgsl/lexer.rs

@@ -347,12 +347,26 @@ impl<'a> Lexer<'a> {
         Ok(pair)
     }
 
-    pub(super) fn next_format_generic(&mut self) -> Result<crate::StorageFormat, Error<'a>> {
+    // TODO relocate storage texture specifics
+    pub(super) fn next_format_generic(
+        &mut self,
+    ) -> Result<(crate::StorageFormat, crate::StorageAccess), Error<'a>> {
         self.expect(Token::Paren('<'))?;
         let (ident, ident_span) = self.next_ident_with_span()?;
         let format = conv::map_storage_format(ident, ident_span)?;
+        let access = if self.skip(Token::Separator(',')) {
+            let (raw, span) = self.next_ident_with_span()?;
+            match raw {
+                "read" => crate::StorageAccess::LOAD,
+                "write" => crate::StorageAccess::STORE,
+                "read_write" => crate::StorageAccess::all(),
+                _ => return Err(Error::UnknownAccess(span)),
+            }
+        } else {
+            crate::StorageAccess::LOAD
+        };
         self.expect(Token::Paren('>'))?;
-        Ok(format)
+        Ok((format, access))
     }
 
     pub(super) fn open_arguments(&mut self) -> Result<(), Error<'a>> {
@@ -455,5 +469,23 @@ fn test_variable_decl() {
             Token::Paren('>'),
             Token::Separator(';'),
         ],
-    )
+    );
+    sub_test(
+        "var<storage,read_write> buffer: array<u32>;",
+        &[
+            Token::Word("var"),
+            Token::Paren('<'),
+            Token::Word("storage"),
+            Token::Separator(','),
+            Token::Word("read_write"),
+            Token::Paren('>'),
+            Token::Word("buffer"),
+            Token::Separator(':'),
+            Token::Word("array"),
+            Token::Paren('<'),
+            Token::Word("u32"),
+            Token::Paren('>'),
+            Token::Separator(';'),
+        ],
+    );
 }

src/front/wgsl/mod.rs

@@ -467,8 +467,12 @@ impl crate::TypeInner {
                         format!("<{}>", element_type)
                     }
                     crate::ImageClass::Depth { multi: _ } => String::new(),
-                    crate::ImageClass::Storage(format) => {
-                        format!("<{}>", format.to_wgsl())
+                    crate::ImageClass::Storage { format, access } => {
+                        if access.contains(crate::StorageAccess::STORE) {
+                            format!("<{},write>", format.to_wgsl())
+                        } else {
+                            format!("<{}>", format.to_wgsl())
+                        }
                     }
                 };
 
@@ -529,7 +533,9 @@ mod type_inner_tests {
 
         let ptr = crate::TypeInner::Pointer {
             base: mytype2,
-            class: crate::StorageClass::Storage,
+            class: crate::StorageClass::Storage {
+                access: crate::StorageAccess::default(),
+            },
         };
         assert_eq!(ptr.to_wgsl(&types, &constants), "*MyType2");
 
@@ -836,7 +842,6 @@ impl Composition {
 #[derive(Default)]
 struct TypeAttributes {
     stride: Option<NonZeroU32>,
-    access: crate::StorageAccess,
 }
 
 #[derive(Clone, Debug, PartialEq)]
@@ -927,7 +932,6 @@ struct ParsedVariable<'a> {
     name: &'a str,
     class: Option<crate::StorageClass>,
     ty: Handle<crate::Type>,
-    access: crate::StorageAccess,
     init: Option<Handle<crate::Constant>>,
 }
 
@@ -1352,7 +1356,7 @@ impl Parser {
                         None
                     };
                     let index = match class {
-                        crate::ImageClass::Storage(_) => None,
+                        crate::ImageClass::Storage { .. } => None,
                         // it's the MSAA index for multi-sampled, and LOD for the others
                         crate::ImageClass::Sampled { .. } | crate::ImageClass::Depth { .. } => {
                             lexer.expect(Token::Separator(','))?;
@@ -1997,14 +2001,32 @@ impl Parser {
     ) -> Result<ParsedVariable<'a>, Error<'a>> {
         self.scopes.push(Scope::VariableDecl);
         let mut class = None;
+
         if lexer.skip(Token::Paren('<')) {
             let (class_str, span) = lexer.next_ident_with_span()?;
-            class = Some(conv::map_storage_class(class_str, span)?);
+            class = Some(match class_str {
+                "storage" => {
+                    let access = if lexer.skip(Token::Separator(',')) {
+                        let (ident, span) = lexer.next_ident_with_span()?;
+                        match ident {
+                            "read" => crate::StorageAccess::LOAD,
+                            "write" => crate::StorageAccess::STORE,
+                            "read_write" => crate::StorageAccess::all(),
+                            _ => return Err(Error::UnknownAccess(span)),
+                        }
+                    } else {
+                        // defaulting to `read`
+                        crate::StorageAccess::LOAD
+                    };
+                    crate::StorageClass::Storage { access }
+                }
+                _ => conv::map_storage_class(class_str, span)?,
+            });
             lexer.expect(Token::Paren('>'))?;
         }
         let name = lexer.next_ident()?;
         lexer.expect(Token::Separator(':'))?;
-        let (ty, access) = self.parse_type_decl(lexer, None, type_arena, const_arena)?;
+        let (ty, _access) = self.parse_type_decl(lexer, None, type_arena, const_arena)?;
 
         let init = if lexer.skip(Token::Operation('=')) {
             let handle = self.parse_const_expression(lexer, type_arena, const_arena)?;
@@ -2018,7 +2040,6 @@ impl Parser {
             name,
             class,
             ty,
-            access,
             init,
         })
     }
@@ -2363,43 +2384,43 @@ impl Parser {
                 class: crate::ImageClass::Depth { multi: true },
             },
             "texture_storage_1d" => {
-                let format = lexer.next_format_generic()?;
+                let (format, access) = lexer.next_format_generic()?;
                 crate::TypeInner::Image {
                     dim: crate::ImageDimension::D1,
                     arrayed: false,
-                    class: crate::ImageClass::Storage(format),
+                    class: crate::ImageClass::Storage { format, access },
                 }
             }
             "texture_storage_1d_array" => {
-                let format = lexer.next_format_generic()?;
+                let (format, access) = lexer.next_format_generic()?;
                 crate::TypeInner::Image {
                     dim: crate::ImageDimension::D1,
                     arrayed: true,
-                    class: crate::ImageClass::Storage(format),
+                    class: crate::ImageClass::Storage { format, access },
                 }
             }
             "texture_storage_2d" => {
-                let format = lexer.next_format_generic()?;
+                let (format, access) = lexer.next_format_generic()?;
                 crate::TypeInner::Image {
                     dim: crate::ImageDimension::D2,
                     arrayed: false,
-                    class: crate::ImageClass::Storage(format),
+                    class: crate::ImageClass::Storage { format, access },
                 }
             }
             "texture_storage_2d_array" => {
-                let format = lexer.next_format_generic()?;
+                let (format, access) = lexer.next_format_generic()?;
                 crate::TypeInner::Image {
                     dim: crate::ImageDimension::D2,
                     arrayed: true,
-                    class: crate::ImageClass::Storage(format),
+                    class: crate::ImageClass::Storage { format, access },
                 }
             }
             "texture_storage_3d" => {
-                let format = lexer.next_format_generic()?;
+                let (format, access) = lexer.next_format_generic()?;
                 crate::TypeInner::Image {
                     dim: crate::ImageDimension::D3,
                     arrayed: false,
-                    class: crate::ImageClass::Storage(format),
+                    class: crate::ImageClass::Storage { format, access },
                 }
             }
             _ => return Ok(None),
@@ -2459,17 +2480,6 @@ impl Parser {
             self.scopes.push(Scope::Attribute);
             loop {
                 match lexer.next() {
-                    (Token::Word("access"), _) => {
-                        lexer.expect(Token::Paren('('))?;
-                        let (ident, span) = lexer.next_ident_with_span()?;
-                        attribute.access = match ident {
-                            "read" => crate::StorageAccess::LOAD,
-                            "write" => crate::StorageAccess::STORE,
-                            "read_write" => crate::StorageAccess::all(),
-                            _ => return Err(Error::UnknownAccess(span)),
-                        };
-                        lexer.expect(Token::Paren(')'))?;
-                    }
                     (Token::Word("stride"), _) => {
                         lexer.expect(Token::Paren('('))?;
                         let (stride, span) = lexer.capture_span(Lexer::next_uint_literal)?;
@@ -2484,7 +2494,7 @@ impl Parser {
             self.scopes.pop();
         }
 
-        let storage_access = attribute.access;
+        let storage_access = crate::StorageAccess::default();
         let (name, name_span) = lexer.next_ident_with_span()?;
         let handle = self.parse_type_decl_name(
             lexer,
@@ -3271,14 +3281,12 @@ impl Parser {
                     Some(c) => c,
                     None => match module.types[pvar.ty].inner {
                         crate::TypeInner::Struct { .. } if binding.is_some() => {
-                            if pvar.access.is_empty() {
-                                crate::StorageClass::Uniform
-                            } else {
-                                crate::StorageClass::Storage
-                            }
+                            crate::StorageClass::Uniform
                         }
                         crate::TypeInner::Array { .. } if binding.is_some() => {
-                            crate::StorageClass::Storage
+                            crate::StorageClass::Storage {
+                                access: crate::StorageAccess::LOAD,
+                            }
                         }
                         crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {
                             crate::StorageClass::Handle
@@ -3292,7 +3300,6 @@ impl Parser {
                     binding: binding.take(),
                     ty: pvar.ty,
                     init: pvar.init,
-                    storage_access: pvar.access,
                 });
                 lookup_global_expression
                     .insert(pvar.name, crate::Expression::GlobalVariable(var_handle));

src/front/wgsl/tests.rs

@@ -21,8 +21,8 @@ fn parse_types() {
     parse_str("var t: texture_2d<f32>;").unwrap();
     parse_str("var t: texture_cube_array<i32>;").unwrap();
     parse_str("var t: texture_multisampled_2d<u32>;").unwrap();
-    parse_str("var t: [[access(write)]] texture_storage_1d<rgba8uint>;").unwrap();
-    parse_str("var t: [[access(read)]] texture_storage_3d<r32float>;").unwrap();
+    parse_str("var t: texture_storage_1d<rgba8uint,write>;").unwrap();
+    parse_str("var t: texture_storage_3d<r32float>;").unwrap();
 }
 
 #[test]
@@ -78,7 +78,7 @@ fn parse_struct() {
             [[size(32), align(8)]] z: vec3<f32>;
         };
         struct Empty {};
-        var s: [[access(read_write)]] Foo;
+        var<storage,read_write> s: Foo;
     ",
     )
     .unwrap();
@@ -219,7 +219,7 @@ fn parse_texture_load() {
     .unwrap();
     parse_str(
         "
-        var t: [[access(read)]] texture_storage_1d_array<r32float>;
+        var t: texture_storage_1d_array<r32float>;
         fn foo() {
             let r: vec4<f32> = textureLoad(t, 10, 2);
         }
@@ -232,7 +232,7 @@ fn parse_texture_load() {
 fn parse_texture_store() {
     parse_str(
         "
-        var t: [[access(write)]] texture_storage_2d<rgba8unorm>;
+        var t: texture_storage_2d<rgba8unorm,write>;
         fn foo() {
             textureStore(t, vec2<i32>(10, 20), vec4<f32>(0.0, 1.0, 2.0, 3.0));
         }
@@ -317,10 +317,10 @@ fn parse_array_length() {
         }; // this is used as both input and output for convenience
 
         [[group(0), binding(0)]]
-        var<storage> foo: [[access(read_write)]] Foo;
+        var<storage> foo: Foo;
 
         [[group(0), binding(1)]]
-        var<storage> bar: [[access(read)]] array<u32>;
+        var<storage> bar: array<u32>;
 
         fn foo() {
             var x: u32 = arrayLength(foo.data);
@@ -330,3 +330,35 @@ fn parse_array_length() {
     )
     .unwrap();
 }
+
+#[test]
+fn parse_storage_buffers() {
+    parse_str(
+        "
+        [[group(0), binding(0)]]
+        var<storage> foo: array<u32>;
+        ",
+    )
+    .unwrap();
+    parse_str(
+        "
+        [[group(0), binding(0)]]
+        var<storage,read> foo: array<u32>;
+        ",
+    )
+    .unwrap();
+    parse_str(
+        "
+        [[group(0), binding(0)]]
+        var<storage,write> foo: array<u32>;
+        ",
+    )
+    .unwrap();
+    parse_str(
+        "
+        [[group(0), binding(0)]]
+        var<storage,read_write> foo: array<u32>;
+        ",
+    )
+    .unwrap();
+}

src/lib.rs

@@ -245,7 +245,7 @@ pub enum StorageClass {
     /// Uniform buffer data.
     Uniform,
     /// Storage buffer data, potentially mutable.
-    Storage,
+    Storage { access: StorageAccess },
     /// Opaque handles, such as samplers and images.
     Handle,
     /// Push constants.
@@ -473,7 +473,10 @@ pub enum ImageClass {
         multi: bool,
     },
     /// Storage image.
-    Storage(StorageFormat),
+    Storage {
+        format: StorageFormat,
+        access: StorageAccess,
+    },
 }
 
 /// A data type declared in the module.
@@ -682,8 +685,6 @@ pub struct GlobalVariable {
     pub ty: Handle<Type>,
     /// Initial value for this variable.
     pub init: Option<Handle<Constant>>,
-    /// Access bit for storage types of images and buffers.
-    pub storage_access: StorageAccess,
 }
 
 /// Variable defined at function level.

src/proc/typifier.rs

@@ -377,7 +377,7 @@ impl<'a> ResolveContext<'a> {
                         width: 4,
                         size: crate::VectorSize::Quad,
                     },
-                    crate::ImageClass::Storage(format) => Ti::Vector {
+                    crate::ImageClass::Storage { format, .. } => Ti::Vector {
                         kind: format.into(),
                         width: 4,
                         size: crate::VectorSize::Quad,

src/valid/analyzer.rs

@@ -434,9 +434,8 @@ impl FunctionInfo {
                     // uniform data
                     Sc::Uniform | Sc::PushConstant => true,
                     // storage data is only uniform when read-only
-                    Sc::Handle | Sc::Storage => {
-                        !var.storage_access.contains(crate::StorageAccess::STORE)
-                    }
+                    Sc::Storage { access } => !access.contains(crate::StorageAccess::STORE),
+                    Sc::Handle => false,
                 };
                 Uniformity {
                     non_uniform_result: if uniform { None } else { Some(handle) },
@@ -865,7 +864,6 @@ fn uniform_control_flow() {
         ty,
         class: crate::StorageClass::Handle,
         binding: None,
-        storage_access: crate::StorageAccess::STORE,
     });
     let uniform_global = global_var_arena.append(crate::GlobalVariable {
         name: None,
@@ -873,7 +871,6 @@ fn uniform_control_flow() {
         ty,
         binding: None,
         class: crate::StorageClass::Uniform,
-        storage_access: crate::StorageAccess::empty(),
     });
 
     let mut expressions = Arena::new();

src/valid/function.rs

@@ -475,11 +475,13 @@ impl super::Validator {
                                 }
                             }
                             match class {
-                                crate::ImageClass::Storage(format) => crate::TypeInner::Vector {
-                                    kind: format.into(),
-                                    size: crate::VectorSize::Quad,
-                                    width: 4,
-                                },
+                                crate::ImageClass::Storage { format, .. } => {
+                                    crate::TypeInner::Vector {
+                                        kind: format.into(),
+                                        size: crate::VectorSize::Quad,
+                                        width: 4,
+                                    }
+                                }
                                 _ => {
                                     return Err(FunctionError::InvalidImageStore(
                                         ExpressionError::InvalidImageClass(class),

src/valid/interface.rs

@@ -15,11 +15,6 @@ pub enum GlobalVariableError {
     InvalidUsage,
     #[error("Type isn't compatible with the storage class")]
     InvalidType,
-    #[error("Storage access {seen:?} exceeds the allowed {allowed:?}")]
-    InvalidStorageAccess {
-        allowed: crate::StorageAccess,
-        seen: crate::StorageAccess,
-    },
     #[error("Type flags {seen:?} do not meet the required {required:?}")]
     MissingTypeFlags {
         required: TypeFlags,
@@ -322,16 +317,15 @@ impl super::Validator {
         log::debug!("var {:?}", var);
         let type_info = &self.types[var.ty.index()];
 
-        let (allowed_storage_access, required_type_flags, is_resource) = match var.class {
+        let (required_type_flags, is_resource) = match var.class {
             crate::StorageClass::Function => return Err(GlobalVariableError::InvalidUsage),
-            crate::StorageClass::Storage => {
+            crate::StorageClass::Storage { .. } => {
                 if let Err((ty_handle, disalignment)) = type_info.storage_layout {
                     if self.flags.contains(ValidationFlags::STRUCT_LAYOUTS) {
                         return Err(GlobalVariableError::Alignment(ty_handle, disalignment));
                     }
                 }
                 (
-                    crate::StorageAccess::all(),
                     TypeFlags::DATA | TypeFlags::HOST_SHARED | TypeFlags::TOP_LEVEL,
                     true,
                 )
@@ -343,7 +337,6 @@ impl super::Validator {
                     }
                 }
                 (
-                    crate::StorageAccess::empty(),
                     TypeFlags::DATA
                         | TypeFlags::SIZED
                         | TypeFlags::HOST_SHARED
@@ -352,23 +345,15 @@ impl super::Validator {
                 )
             }
             crate::StorageClass::Handle => {
-                let access = match types[var.ty].inner {
-                    crate::TypeInner::Image {
-                        class: crate::ImageClass::Storage(_),
-                        ..
-                    } => crate::StorageAccess::all(),
-                    crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {
-                        crate::StorageAccess::empty()
-                    }
+                match types[var.ty].inner {
+                    crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {}
                     _ => return Err(GlobalVariableError::InvalidType),
                 };
-                (access, TypeFlags::empty(), true)
+                (TypeFlags::empty(), true)
+            }
+            crate::StorageClass::Private | crate::StorageClass::WorkGroup => {
+                (TypeFlags::DATA | TypeFlags::SIZED, false)
             }
-            crate::StorageClass::Private | crate::StorageClass::WorkGroup => (
-                crate::StorageAccess::empty(),
-                TypeFlags::DATA | TypeFlags::SIZED,
-                false,
-            ),
             crate::StorageClass::PushConstant => {
                 if !self.capabilities.contains(Capabilities::PUSH_CONSTANT) {
                     return Err(GlobalVariableError::UnsupportedCapability(
@@ -376,20 +361,12 @@ impl super::Validator {
                     ));
                 }
                 (
-                    crate::StorageAccess::LOAD,
                     TypeFlags::DATA | TypeFlags::HOST_SHARED | TypeFlags::SIZED,
                     false,
                 )
             }
         };
 
-        if !allowed_storage_access.contains(var.storage_access) {
-            return Err(GlobalVariableError::InvalidStorageAccess {
-                seen: var.storage_access,
-                allowed: allowed_storage_access,
-            });
-        }
-
         if !type_info.flags.contains(required_type_flags) {
             return Err(GlobalVariableError::MissingTypeFlags {
                 seen: type_info.flags,
@@ -481,12 +458,12 @@ impl super::Validator {
             let allowed_usage = match var.class {
                 crate::StorageClass::Function => unreachable!(),
                 crate::StorageClass::Uniform => GlobalUse::READ | GlobalUse::QUERY,
-                crate::StorageClass::Storage => storage_usage(var.storage_access),
+                crate::StorageClass::Storage { access } => storage_usage(access),
                 crate::StorageClass::Handle => match module.types[var.ty].inner {
                     crate::TypeInner::Image {
-                        class: crate::ImageClass::Storage(_),
+                        class: crate::ImageClass::Storage { access, .. },
                         ..
-                    } => storage_usage(var.storage_access),
+                    } => storage_usage(access),
                     _ => GlobalUse::READ | GlobalUse::QUERY,
                 },
                 crate::StorageClass::Private | crate::StorageClass::WorkGroup => GlobalUse::all(),

tests/in/access.wgsl

@@ -8,7 +8,7 @@ struct Bar {
 };
 
 [[group(0), binding(0)]]
-var<storage> bar: [[access(read_write)]] Bar;
+var<storage,read_write> bar: Bar;
 
 [[stage(vertex)]]
 fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {

tests/in/boids.wgsl

@@ -22,8 +22,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> 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)]]

tests/in/collatz.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

tests/in/image.wgsl

@@ -5,9 +5,9 @@ var image_multisampled_src: texture_multisampled_2d<u32>;
 [[group(0), binding(4)]]
 var image_depth_multisampled_src: texture_depth_multisampled_2d;
 [[group(0), binding(1)]]
-var image_storage_src: [[access(read)]] texture_storage_2d<rgba8uint>;
+var image_storage_src: texture_storage_2d<rgba8uint>;
 [[group(0), binding(2)]]
-var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
+var image_dst: texture_storage_1d<r32uint,write>;
 
 [[stage(compute), workgroup_size(16)]]
 fn main(

tests/in/shadow.wgsl

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

tests/out/analysis/collatz.info.ron

@@ -32,7 +32,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 3,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -372,7 +376,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 3,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -397,7 +405,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 2,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -425,7 +437,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 1,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -439,7 +455,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 2,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -467,7 +487,11 @@
                     assignable_global: Some(1),
                     ty: Value(Pointer(
                         base: 1,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (

tests/out/analysis/shadow.info.ron

@@ -8,7 +8,7 @@
                 bits: 7,
             ),
             uniformity: (
-                non_uniform_result: Some(48),
+                non_uniform_result: Some(4),
                 requirements: (
                     bits: 0,
                 ),
@@ -88,7 +88,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(4),
                         requirements: (
                             bits: 0,
                         ),
@@ -99,7 +99,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(5),
                         requirements: (
                             bits: 0,
                         ),
@@ -110,7 +110,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -119,7 +119,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 21,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -1047,7 +1051,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(4),
                         requirements: (
                             bits: 0,
                         ),
@@ -1070,7 +1074,7 @@
                 bits: 7,
             ),
             uniformity: (
-                non_uniform_result: Some(48),
+                non_uniform_result: Some(4),
                 requirements: (
                     bits: 0,
                 ),
@@ -1150,7 +1154,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(4),
                         requirements: (
                             bits: 0,
                         ),
@@ -1161,7 +1165,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(5),
                         requirements: (
                             bits: 0,
                         ),
@@ -1172,7 +1176,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1181,7 +1185,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 21,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -1876,7 +1884,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1885,7 +1893,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -1901,7 +1913,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1910,12 +1922,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1924,12 +1940,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 18,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1951,7 +1971,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -1966,7 +1986,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(4),
                         requirements: (
                             bits: 0,
                         ),
@@ -1999,7 +2019,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2008,7 +2028,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2024,7 +2048,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2033,12 +2057,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2047,12 +2075,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2063,12 +2095,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2082,7 +2118,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2091,7 +2127,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2107,7 +2147,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2116,12 +2156,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2130,12 +2174,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2146,12 +2194,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2165,7 +2217,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2174,7 +2226,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2190,7 +2246,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2199,12 +2255,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2213,12 +2273,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2229,12 +2293,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2248,7 +2316,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2360,7 +2428,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2371,7 +2439,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2410,7 +2478,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(4),
                         requirements: (
                             bits: 0,
                         ),
@@ -2424,7 +2492,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2433,7 +2501,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2449,7 +2521,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2458,12 +2530,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2472,12 +2548,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2488,12 +2568,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2507,7 +2591,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2516,7 +2600,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2532,7 +2620,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2541,12 +2629,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2555,12 +2647,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2571,12 +2667,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2590,7 +2690,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2599,7 +2699,11 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 20,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
@@ -2615,7 +2719,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2624,12 +2728,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 19,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2638,12 +2746,16 @@
                     assignable_global: Some(4),
                     ty: Value(Pointer(
                         base: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2654,12 +2766,16 @@
                         size: None,
                         kind: Float,
                         width: 4,
-                        class: Storage,
+                        class: Storage(
+                            access: (
+                                bits: 3,
+                            ),
+                        ),
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2673,7 +2789,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2684,7 +2800,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(48),
+                        non_uniform_result: Some(6),
                         requirements: (
                             bits: 0,
                         ),
@@ -2761,7 +2877,7 @@
                 bits: 7,
             ),
             uniformity: (
-                non_uniform_result: Some(48),
+                non_uniform_result: Some(4),
                 requirements: (
                     bits: 0,
                 ),

tests/out/dot/quad.dot

@@ -35,9 +35,9 @@ digraph Module {
 		label="Fragment/'main'"
 		node [ style=filled ]
 		ep1_e0 [ fillcolor="#ffffb3" label="[1] Constant" ]
-		ep1_e1 [ fillcolor="#ffffb3" label="[2] Global" ]
+		ep1_e1 [ color="#ffffb3" label="[2] Global" ]
 		g1 -> ep1_e1 [fillcolor=gray]
-		ep1_e2 [ fillcolor="#ffffb3" label="[3] Global" ]
+		ep1_e2 [ color="#ffffb3" label="[3] Global" ]
 		g0 -> ep1_e2 [fillcolor=gray]
 		ep1_e3 [ color="#8dd3c7" label="[4] Argument[0]" ]
 		ep1_e4 [ color="#80b1d3" label="[5] ImageSample" ]
@@ -86,9 +86,9 @@ digraph Module {
 		label="Fragment/'fs_extra'"
 		node [ style=filled ]
 		ep2_e0 [ fillcolor="#ffffb3" label="[1] Constant" ]
-		ep2_e1 [ fillcolor="#ffffb3" label="[2] Global" ]
+		ep2_e1 [ color="#ffffb3" label="[2] Global" ]
 		g1 -> ep2_e1 [fillcolor=gray]
-		ep2_e2 [ fillcolor="#ffffb3" label="[3] Global" ]
+		ep2_e2 [ color="#ffffb3" label="[3] Global" ]
 		g0 -> ep2_e2 [fillcolor=gray]
 		ep2_e3 [ fillcolor="#ffffb3" label="[4] Constant" ]
 		ep2_e4 [ fillcolor="#ffffb3" label="[5] Constant" ]

tests/out/hlsl/image.hlsl

@@ -2,7 +2,7 @@
 Texture2D<uint4> image_mipmapped_src : register(t0);
 Texture2DMS<uint4> image_multisampled_src : register(t3);
 Texture2DMS<float> image_depth_multisampled_src : register(t4);
-Texture2D<uint4> image_storage_src : register(t1);
+RWTexture2D<uint4> image_storage_src : register(u1);
 RWTexture1D<uint4> image_dst : register(u2);
 Texture1D<float4> image_1d : register(t0);
 Texture2D<float4> image_2d : register(t1);
@@ -19,22 +19,22 @@ struct ComputeInput_main {
     uint3 local_id1 : SV_GroupThreadID;
 };
 
-int2 NagaDimensions2D(Texture2D<uint4> texture)
+int2 NagaRWDimensions2D(RWTexture2D<uint4> texture)
 {
     uint4 ret;
-    texture.GetDimensions(0, ret.x, ret.y, ret.z);
+    texture.GetDimensions(ret.x, ret.y);
     return ret.xy;
 }
 
 [numthreads(16, 1, 1)]
 void main(ComputeInput_main computeinput_main)
 {
-    int2 dim = NagaDimensions2D(image_storage_src);
+    int2 dim = NagaRWDimensions2D(image_storage_src);
     int2 itc = ((dim * int2(computeinput_main.local_id1.xy)) % int2(10, 20));
     uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(computeinput_main.local_id1.z)));
     uint4 value2_ = image_multisampled_src.Load(itc, int(computeinput_main.local_id1.z));
     float value3_ = image_depth_multisampled_src.Load(itc, int(computeinput_main.local_id1.z)).x;
-    uint4 value4_ = image_storage_src.Load(int3(itc, 0));
+    uint4 value4_ = image_storage_src.Load(itc);
     image_dst[itc.x] = (((value1_ + value2_) + uint4(uint(value3_).xxxx)) + value4_);
     return;
 }
@@ -42,7 +42,7 @@ void main(ComputeInput_main computeinput_main)
 int NagaDimensions1D(Texture1D<float4> texture)
 {
     uint4 ret;
-    texture.GetDimensions(ret.x);
+    texture.GetDimensions(0, ret.x, ret.y);
     return ret.x;
 }
 

tests/out/ir/collatz.ron

@@ -76,16 +76,17 @@
     global_variables: [
         (
             name: Some("v_indices"),
-            class: Storage,
+            class: Storage(
+                access: (
+                    bits: 3,
+                ),
+            ),
             binding: Some((
                 group: 0,
                 binding: 0,
             )),
             ty: 3,
             init: None,
-            storage_access: (
-                bits: 3,
-            ),
         ),
     ],
     functions: [

tests/out/ir/shadow.ron

@@ -197,7 +197,11 @@
             name: None,
             inner: Pointer(
                 base: 21,
-                class: Storage,
+                class: Storage(
+                    access: (
+                        bits: 0,
+                    ),
+                ),
             ),
         ),
         (
@@ -776,9 +780,6 @@
             )),
             ty: 56,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
         (
             name: Some("sampler_shadow"),
@@ -789,9 +790,6 @@
             )),
             ty: 57,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
         (
             name: Some("u_globals"),
@@ -802,22 +800,20 @@
             )),
             ty: 14,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
         (
             name: Some("s_lights"),
-            class: Storage,
+            class: Storage(
+                access: (
+                    bits: 3,
+                ),
+            ),
             binding: Some((
                 group: 0,
                 binding: 1,
             )),
             ty: 21,
             init: None,
-            storage_access: (
-                bits: 1,
-            ),
         ),
         (
             name: Some("in_position_fs"),
@@ -825,9 +821,6 @@
             binding: None,
             ty: 4,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
         (
             name: Some("in_normal_fs"),
@@ -835,9 +828,6 @@
             binding: None,
             ty: 2,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
         (
             name: Some("out_color_fs"),
@@ -845,9 +835,6 @@
             binding: None,
             ty: 4,
             init: None,
-            storage_access: (
-                bits: 0,
-            ),
         ),
     ],
     functions: [

tests/out/spv/bounds-check-zero.spvasm

@@ -14,6 +14,7 @@ OpMemberDecorate %9 1 Offset 48
 OpMemberDecorate %9 2 Offset 64
 OpMemberDecorate %9 2 ColMajor
 OpMemberDecorate %9 2 MatrixStride 16
+OpDecorate %10 NonWritable
 OpDecorate %10 DescriptorSet 0
 OpDecorate %10 Binding 0
 %2 = OpTypeVoid

tests/out/spv/pointer-access.spvasm

@@ -12,6 +12,8 @@ OpDecorate %12 ArrayStride 4
 OpDecorate %14 ArrayStride 4
 OpDecorate %15 Block
 OpMemberDecorate %15 0 Offset 0
+OpDecorate %17 NonReadable
+OpDecorate %17 NonWritable
 OpDecorate %17 DescriptorSet 0
 OpDecorate %17 Binding 0
 %2 = OpTypeVoid

tests/out/wgsl/246-collatz-comp.wgsl

@@ -4,7 +4,7 @@ struct PrimeIndices {
 };
 
 [[group(0), binding(0)]]
-var<storage> global: [[access(read_write)]] PrimeIndices;
+var<storage,read_write> global: PrimeIndices;
 var<private> gl_GlobalInvocationID: vec3<u32>;
 
 fn collatz_iterations(n: u32) -> u32 {

tests/out/wgsl/access.wgsl

@@ -6,7 +6,7 @@ struct Bar {
 };
 
 [[group(0), binding(0)]]
-var<storage> bar: [[access(read_write)]] Bar;
+var<storage,read_write> bar: Bar;
 
 [[stage(vertex)]]
 fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {

tests/out/wgsl/boids.wgsl

@@ -24,9 +24,9 @@ let NUM_PARTICLES: u32 = 1500u;
 [[group(0), binding(0)]]
 var<uniform> params: SimParams;
 [[group(0), binding(1)]]
-var<storage> particlesSrc: [[access(read)]] Particles;
+var<storage> particlesSrc: Particles;
 [[group(0), binding(2)]]
-var<storage> particlesDst: [[access(read_write)]] Particles;
+var<storage,read_write> particlesDst: Particles;
 
 [[stage(compute), workgroup_size(64, 1, 1)]]
 fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3<u32>) {

tests/out/wgsl/collatz.wgsl

@@ -4,7 +4,7 @@ struct PrimeIndices {
 };
 
 [[group(0), binding(0)]]
-var<storage> v_indices: [[access(read_write)]] PrimeIndices;
+var<storage,read_write> v_indices: PrimeIndices;
 
 fn collatz_iterations(n_base: u32) -> u32 {
     var n: u32;

tests/out/wgsl/image.wgsl

@@ -5,9 +5,9 @@ var image_multisampled_src: texture_multisampled_2d<u32>;
 [[group(0), binding(4)]]
 var image_depth_multisampled_src: texture_depth_multisampled_2d;
 [[group(0), binding(1)]]
-var image_storage_src: [[access(read)]] texture_storage_2d<rgba8uint>;
+var image_storage_src: texture_storage_2d<rgba8uint>;
 [[group(0), binding(2)]]
-var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
+var image_dst: texture_storage_1d<r32uint,write>;
 [[group(0), binding(0)]]
 var image_1d: texture_1d<f32>;
 [[group(0), binding(1)]]

tests/out/wgsl/shadow.wgsl

@@ -20,7 +20,7 @@ let c_max_lights: u32 = 10u;
 [[group(0), binding(0)]]
 var<uniform> u_globals: Globals;
 [[group(0), binding(1)]]
-var<storage> s_lights: [[access(read)]] Lights;
+var<storage> s_lights: Lights;
 [[group(0), binding(2)]]
 var t_shadow: texture_depth_2d_array;
 [[group(0), binding(3)]]

tests/wgsl-errors.rs

@@ -261,13 +261,13 @@ fn unknown_built_in() {
 fn unknown_access() {
     check(
         r#"
-            var<storage> x: [[access(unknown_access)]] array<u32>;
+            var<storage,unknown_access> x: array<u32>;
         "#,
         r#"error: unknown access: 'unknown_access'
-  ┌─ wgsl:2:38
+  ┌─ wgsl:2:25
   │
-2 │             var<storage> x: [[access(unknown_access)]] array<u32>;
-  │                                      ^^^^^^^^^^^^^^ unknown access
+2 │             var<storage,unknown_access> x: array<u32>;
+  │                         ^^^^^^^^^^^^^^ unknown access
 
 "#,
     );
@@ -345,13 +345,13 @@ fn unknown_type() {
 fn unknown_storage_format() {
     check(
         r#"
-            let storage: [[access(read)]] texture_storage_1d<rgba>;
+            let storage: texture_storage_1d<rgba>;
         "#,
         r#"error: unknown storage format: 'rgba'
-  ┌─ wgsl:2:62
+  ┌─ wgsl:2:45
   │
-2 │             let storage: [[access(read)]] texture_storage_1d<rgba>;
-  │                                                              ^^^^ unknown storage format
+2 │             let storage: texture_storage_1d<rgba>;
+  │                                             ^^^^ unknown storage format
 
 "#,
     );