diff --git a/CHANGELOG.md b/CHANGELOG.md
index 23741c882a..6c7a8dbc7c 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -118,6 +118,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 
 #### General
 
+- Align Storage Access enums to the webgpu spec. By @atlv24 in [#6642](https://github.com/gfx-rs/wgpu/pull/6642)
 - Make `Surface::as_hal` take an immutable reference to the surface. By @jerzywilczek in [#9999](https://github.com/gfx-rs/wgpu/pull/9999)
 - Add actual sample type to `CreateBindGroupError::InvalidTextureSampleType` error message. By @ErichDonGubler in [#6530](https://github.com/gfx-rs/wgpu/pull/6530).
 - Improve binding error to give a clearer message when there is a mismatch between resource binding as it is in the shader and as it is in the binding layout. By @eliemichel in [#6553](https://github.com/gfx-rs/wgpu/pull/6553).
diff --git a/examples/src/ray_traced_triangle/mod.rs b/examples/src/ray_traced_triangle/mod.rs
index d508e6113e..aeda991fa2 100644
--- a/examples/src/ray_traced_triangle/mod.rs
+++ b/examples/src/ray_traced_triangle/mod.rs
@@ -30,6 +30,7 @@ impl crate::framework::Example for Example {
     fn required_features() -> wgpu::Features {
         wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
             | wgpu::Features::EXPERIMENTAL_RAY_QUERY
+            | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES
     }
 
     fn required_limits() -> wgpu::Limits {
diff --git a/player/tests/data/zero-init-texture-binding.ron b/player/tests/data/zero-init-texture-binding.ron
index 48415f43c5..2368a08023 100644
--- a/player/tests/data/zero-init-texture-binding.ron
+++ b/player/tests/data/zero-init-texture-binding.ron
@@ -1,5 +1,5 @@
 (
-    features: [],
+    features: ["TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES"],
     expectations: [
         (
             name: "Sampled Texture",
diff --git a/tests/tests/bgra8unorm_storage.rs b/tests/tests/bgra8unorm_storage.rs
index 30babe0d82..b1c9da6abd 100644
--- a/tests/tests/bgra8unorm_storage.rs
+++ b/tests/tests/bgra8unorm_storage.rs
@@ -21,7 +21,10 @@ static BGRA8_UNORM_STORAGE: GpuTestConfiguration = GpuTestConfiguration::new()
                 max_storage_textures_per_shader_stage: 1,
                 ..Default::default()
             })
-            .features(wgpu::Features::BGRA8UNORM_STORAGE),
+            .features(
+                wgpu::Features::BGRA8UNORM_STORAGE
+                    | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
+            ),
     )
     .run_async(|ctx| async move {
         let device = &ctx.device;
diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs
index b0d9b5f68c..c7867ab210 100644
--- a/wgpu-core/src/binding_model.rs
+++ b/wgpu-core/src/binding_model.rs
@@ -185,6 +185,10 @@ pub enum CreateBindGroupError {
     DepthStencilAspect,
     #[error("The adapter does not support read access for storage textures of format {0:?}")]
     StorageReadNotSupported(wgt::TextureFormat),
+    #[error("The adapter does not support write access for storage textures of format {0:?}")]
+    StorageWriteNotSupported(wgt::TextureFormat),
+    #[error("The adapter does not support read-write access for storage textures of format {0:?}")]
+    StorageReadWriteNotSupported(wgt::TextureFormat),
     #[error(transparent)]
     ResourceUsageCompatibility(#[from] ResourceUsageCompatibilityError),
     #[error(transparent)]
diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs
index 89681b8e27..7eaafea469 100644
--- a/wgpu-core/src/command/compute.rs
+++ b/wgpu-core/src/command/compute.rs
@@ -944,7 +944,7 @@ fn dispatch_indirect(
         let src_transition = state
             .intermediate_trackers
             .buffers
-            .set_single(&buffer, hal::BufferUses::STORAGE_READ);
+            .set_single(&buffer, hal::BufferUses::STORAGE_READ_ONLY);
         let src_barrier =
             src_transition.map(|transition| transition.into_hal(&buffer, &state.snatch_guard));
         unsafe {
diff --git a/wgpu-core/src/conv.rs b/wgpu-core/src/conv.rs
index a9f48ff691..b114e9826e 100644
--- a/wgpu-core/src/conv.rs
+++ b/wgpu-core/src/conv.rs
@@ -82,7 +82,7 @@ pub fn map_buffer_usage(usage: wgt::BufferUsages) -> hal::BufferUses {
         usage.contains(wgt::BufferUsages::UNIFORM),
     );
     u.set(
-        hal::BufferUses::STORAGE_READ | hal::BufferUses::STORAGE_READ_WRITE,
+        hal::BufferUses::STORAGE_READ_WRITE,
         usage.contains(wgt::BufferUsages::STORAGE),
     );
     u.set(
@@ -122,7 +122,7 @@ pub fn map_texture_usage(
         usage.contains(wgt::TextureUsages::TEXTURE_BINDING),
     );
     u.set(
-        hal::TextureUses::STORAGE_READ | hal::TextureUses::STORAGE_READ_WRITE,
+        hal::TextureUses::STORAGE_READ_WRITE,
         usage.contains(wgt::TextureUsages::STORAGE_BINDING),
     );
     let is_color = aspect.contains(hal::FormatAspects::COLOR);
@@ -179,7 +179,11 @@ pub fn map_texture_usage_from_hal(uses: hal::TextureUses) -> wgt::TextureUsages
     );
     u.set(
         wgt::TextureUsages::STORAGE_BINDING,
-        uses.contains(hal::TextureUses::STORAGE_READ | hal::TextureUses::STORAGE_READ_WRITE),
+        uses.intersects(
+            hal::TextureUses::STORAGE_READ_ONLY
+                | hal::TextureUses::STORAGE_WRITE_ONLY
+                | hal::TextureUses::STORAGE_READ_WRITE,
+        ),
     );
     u.set(
         wgt::TextureUsages::RENDER_ATTACHMENT,
diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs
index a89bf0b6c2..2143a0f2e4 100644
--- a/wgpu-core/src/device/resource.rs
+++ b/wgpu-core/src/device/resource.rs
@@ -521,7 +521,16 @@ impl Device {
             self.require_downlevel_flags(wgt::DownlevelFlags::INDIRECT_EXECUTION)?;
             // We are going to be reading from it, internally;
             // when validating the content of the buffer
-            usage |= hal::BufferUses::STORAGE_READ | hal::BufferUses::STORAGE_READ_WRITE;
+            if !usage.intersects(
+                hal::BufferUses::STORAGE_READ_ONLY | hal::BufferUses::STORAGE_READ_WRITE,
+            ) {
+                if usage.contains(hal::BufferUses::STORAGE_WRITE_ONLY) {
+                    usage |= hal::BufferUses::STORAGE_READ_WRITE;
+                    usage &= !hal::BufferUses::STORAGE_WRITE_ONLY;
+                } else {
+                    usage |= hal::BufferUses::STORAGE_READ_ONLY;
+                }
+            }
         }
 
         if desc.mapped_at_creation {
@@ -1254,7 +1263,8 @@ impl Device {
                 }
                 TextureViewDimension::D3 => {
                     hal::TextureUses::RESOURCE
-                        | hal::TextureUses::STORAGE_READ
+                        | hal::TextureUses::STORAGE_READ_ONLY
+                        | hal::TextureUses::STORAGE_WRITE_ONLY
                         | hal::TextureUses::STORAGE_READ_WRITE
                 }
                 _ => hal::TextureUses::all(),
@@ -1916,7 +1926,7 @@ impl Device {
             wgt::BufferBindingType::Storage { read_only } => (
                 wgt::BufferUsages::STORAGE,
                 if read_only {
-                    hal::BufferUses::STORAGE_READ
+                    hal::BufferUses::STORAGE_READ_ONLY
                 } else {
                     hal::BufferUses::STORAGE_READ_WRITE
                 },
@@ -2489,24 +2499,31 @@ impl Device {
                 }
 
                 let internal_use = match access {
-                    wgt::StorageTextureAccess::WriteOnly => hal::TextureUses::STORAGE_READ_WRITE,
+                    wgt::StorageTextureAccess::WriteOnly => {
+                        if !view.format_features.flags.intersects(
+                            wgt::TextureFormatFeatureFlags::STORAGE_WRITE_ONLY
+                                | wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE,
+                        ) {
+                            return Err(Error::StorageWriteNotSupported(view.desc.format));
+                        }
+                        hal::TextureUses::STORAGE_WRITE_ONLY
+                    }
                     wgt::StorageTextureAccess::ReadOnly => {
-                        if !view
-                            .format_features
-                            .flags
-                            .contains(wgt::TextureFormatFeatureFlags::STORAGE_WRITE)
-                        {
+                        if !view.format_features.flags.intersects(
+                            wgt::TextureFormatFeatureFlags::STORAGE_READ_ONLY
+                                | wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE,
+                        ) {
                             return Err(Error::StorageReadNotSupported(view.desc.format));
                         }
-                        hal::TextureUses::STORAGE_READ
+                        hal::TextureUses::STORAGE_READ_ONLY
                     }
                     wgt::StorageTextureAccess::ReadWrite => {
                         if !view
                             .format_features
                             .flags
-                            .contains(wgt::TextureFormatFeatureFlags::STORAGE_WRITE)
+                            .contains(wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE)
                         {
-                            return Err(Error::StorageReadNotSupported(view.desc.format));
+                            return Err(Error::StorageReadWriteNotSupported(view.desc.format));
                         }
 
                         hal::TextureUses::STORAGE_READ_WRITE
diff --git a/wgpu-core/src/instance.rs b/wgpu-core/src/instance.rs
index 16278e16f8..2f4365e927 100644
--- a/wgpu-core/src/instance.rs
+++ b/wgpu-core/src/instance.rs
@@ -512,7 +512,9 @@ impl Adapter {
         );
         allowed_usages.set(
             wgt::TextureUsages::STORAGE_BINDING,
-            caps.contains(Tfc::STORAGE_WRITE),
+            caps.intersects(
+                Tfc::STORAGE_WRITE_ONLY | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_READ_WRITE,
+            ),
         );
         allowed_usages.set(
             wgt::TextureUsages::RENDER_ATTACHMENT,
@@ -521,7 +523,15 @@ impl Adapter {
 
         let mut flags = wgt::TextureFormatFeatureFlags::empty();
         flags.set(
-            wgt::TextureFormatFeatureFlags::STORAGE_WRITE,
+            wgt::TextureFormatFeatureFlags::STORAGE_READ_ONLY,
+            caps.contains(Tfc::STORAGE_READ_ONLY),
+        );
+        flags.set(
+            wgt::TextureFormatFeatureFlags::STORAGE_WRITE_ONLY,
+            caps.contains(Tfc::STORAGE_WRITE_ONLY),
+        );
+        flags.set(
+            wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE,
             caps.contains(Tfc::STORAGE_READ_WRITE),
         );
 
diff --git a/wgpu-hal/src/auxil/dxgi/conv.rs b/wgpu-hal/src/auxil/dxgi/conv.rs
index 878dab39e9..ad64f044cc 100644
--- a/wgpu-hal/src/auxil/dxgi/conv.rs
+++ b/wgpu-hal/src/auxil/dxgi/conv.rs
@@ -206,7 +206,8 @@ pub fn map_texture_format_for_resource(
     } else if format.is_depth_stencil_format()
         && usage.intersects(
             crate::TextureUses::RESOURCE
-                | crate::TextureUses::STORAGE_READ
+                | crate::TextureUses::STORAGE_READ_ONLY
+                | crate::TextureUses::STORAGE_WRITE_ONLY
                 | crate::TextureUses::STORAGE_READ_WRITE,
         )
     {
diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs
index b88941c81e..a8940846d6 100644
--- a/wgpu-hal/src/dx12/adapter.rs
+++ b/wgpu-hal/src/dx12/adapter.rs
@@ -670,13 +670,13 @@ impl crate::Adapter for super::Adapter {
         );
         // UAVs use srv_uav_format
         caps.set(
-            Tfc::STORAGE_WRITE,
+            Tfc::STORAGE_WRITE_ONLY,
             data_srv_uav
                 .Support1
                 .contains(Direct3D12::D3D12_FORMAT_SUPPORT1_TYPED_UNORDERED_ACCESS_VIEW),
         );
         caps.set(
-            Tfc::STORAGE_READ_WRITE,
+            Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_WRITE_ONLY,
             data_srv_uav
                 .Support2
                 .contains(Direct3D12::D3D12_FORMAT_SUPPORT2_UAV_TYPED_LOAD),
diff --git a/wgpu-hal/src/dx12/conv.rs b/wgpu-hal/src/dx12/conv.rs
index 8e60f6e064..70ba83fa42 100644
--- a/wgpu-hal/src/dx12/conv.rs
+++ b/wgpu-hal/src/dx12/conv.rs
@@ -128,9 +128,9 @@ pub fn map_buffer_usage_to_state(usage: crate::BufferUses) -> Direct3D12::D3D12_
     if usage.intersects(Bu::VERTEX | Bu::UNIFORM) {
         state |= Direct3D12::D3D12_RESOURCE_STATE_VERTEX_AND_CONSTANT_BUFFER;
     }
-    if usage.intersects(Bu::STORAGE_READ_WRITE) {
+    if usage.intersects(Bu::STORAGE_READ_WRITE | Bu::STORAGE_WRITE_ONLY) {
         state |= Direct3D12::D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
-    } else if usage.intersects(Bu::STORAGE_READ) {
+    } else if usage.intersects(Bu::STORAGE_READ_ONLY) {
         state |= Direct3D12::D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE
             | Direct3D12::D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE;
     }
@@ -168,7 +168,7 @@ pub fn map_texture_usage_to_state(usage: crate::TextureUses) -> Direct3D12::D3D1
     if usage.intersects(Tu::DEPTH_STENCIL_WRITE) {
         state |= Direct3D12::D3D12_RESOURCE_STATE_DEPTH_WRITE;
     }
-    if usage.intersects(Tu::STORAGE_READ | Tu::STORAGE_READ_WRITE) {
+    if usage.intersects(Tu::STORAGE_READ_ONLY | Tu::STORAGE_READ_WRITE) {
         state |= Direct3D12::D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
     }
     state
diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs
index 9e089bbd12..2a7d88592e 100644
--- a/wgpu-hal/src/dx12/device.rs
+++ b/wgpu-hal/src/dx12/device.rs
@@ -577,7 +577,7 @@ impl crate::Device for super::Device {
                 None
             },
             handle_uav: if desc.usage.intersects(
-                crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE,
+                crate::TextureUses::STORAGE_READ_ONLY | crate::TextureUses::STORAGE_READ_WRITE,
             ) {
                 match unsafe { view_desc.to_uav() } {
                     Some(raw_desc) => {
diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs
index 7d2ba2d23d..a95511908e 100644
--- a/wgpu-hal/src/gles/adapter.rs
+++ b/wgpu-hal/src/gles/adapter.rs
@@ -1038,7 +1038,8 @@ impl crate::Adapter for super::Adapter {
         let renderable =
             unfilterable | Tfc::COLOR_ATTACHMENT | sample_count | Tfc::MULTISAMPLE_RESOLVE;
         let filterable_renderable = filterable | renderable | Tfc::COLOR_ATTACHMENT_BLEND;
-        let storage = base | Tfc::STORAGE_WRITE | Tfc::STORAGE_READ_WRITE;
+        let storage =
+            base | Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_WRITE_ONLY;
 
         let feature_fn = |f, caps| {
             if self.shared.features.contains(f) {
diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs
index 03edf30831..892e854b52 100644
--- a/wgpu-hal/src/gles/queue.rs
+++ b/wgpu-hal/src/gles/queue.rs
@@ -1225,7 +1225,9 @@ impl super::Queue {
                     flags |= glow::BUFFER_UPDATE_BARRIER_BIT;
                 }
                 if usage.intersects(
-                    crate::BufferUses::STORAGE_READ | crate::BufferUses::STORAGE_READ_WRITE,
+                    crate::BufferUses::STORAGE_READ_ONLY
+                        | crate::BufferUses::STORAGE_WRITE_ONLY
+                        | crate::BufferUses::STORAGE_READ_WRITE,
                 ) {
                     flags |= glow::SHADER_STORAGE_BARRIER_BIT;
                 }
@@ -1237,7 +1239,9 @@ impl super::Queue {
                     flags |= glow::TEXTURE_FETCH_BARRIER_BIT;
                 }
                 if usage.intersects(
-                    crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE,
+                    crate::TextureUses::STORAGE_READ_ONLY
+                        | crate::TextureUses::STORAGE_WRITE_ONLY
+                        | crate::TextureUses::STORAGE_READ_WRITE,
                 ) {
                     flags |= glow::SHADER_IMAGE_ACCESS_BARRIER_BIT;
                 }
diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs
index 44ee9fb9aa..0a263104e5 100644
--- a/wgpu-hal/src/lib.rs
+++ b/wgpu-hal/src/lib.rs
@@ -1543,8 +1543,10 @@ bitflags!(
         /// Format can be sampled with a min/max reduction sampler.
         const SAMPLED_MINMAX = 1 << 2;
 
+        /// Format can be used as storage with read-only access.
+        const STORAGE_READ_ONLY = 1 << 16;
         /// Format can be used as storage with write-only access.
-        const STORAGE_WRITE = 1 << 3;
+        const STORAGE_WRITE_ONLY = 1 << 3;
         /// Format can be used as storage with read and read/write access.
         const STORAGE_READ_WRITE = 1 << 4;
         /// Format can be used as storage with atomics.
@@ -1675,8 +1677,10 @@ bitflags::bitflags! {
         /// A uniform buffer bound in a bind group.
         const UNIFORM = 1 << 6;
         /// A read-only storage buffer used in a bind group.
-        const STORAGE_READ = 1 << 7;
-        /// A read-write or write-only buffer used in a bind group.
+        const STORAGE_READ_ONLY = 1 << 7;
+        /// A write-only storage buffer used in a bind group.
+        const STORAGE_WRITE_ONLY = 1 << 8;
+        /// A read-write buffer used in a bind group.
         const STORAGE_READ_WRITE = 1 << 8;
         /// The indirect or count buffer in a indirect draw or dispatch.
         const INDIRECT = 1 << 9;
@@ -1688,7 +1692,7 @@ bitflags::bitflags! {
         /// The combination of states that a buffer may be in _at the same time_.
         const INCLUSIVE = Self::MAP_READ.bits() | Self::COPY_SRC.bits() |
             Self::INDEX.bits() | Self::VERTEX.bits() | Self::UNIFORM.bits() |
-            Self::STORAGE_READ.bits() | Self::INDIRECT.bits() | Self::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits() | Self::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits();
+            Self::STORAGE_READ_ONLY.bits() | Self::INDIRECT.bits() | Self::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits() | Self::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits();
         /// The combination of states that a buffer must exclusively be in.
         const EXCLUSIVE = Self::MAP_WRITE.bits() | Self::COPY_DST.bits() | Self::STORAGE_READ_WRITE.bits() | Self::ACCELERATION_STRUCTURE_SCRATCH.bits();
         /// The combination of all usages that the are guaranteed to be be ordered by the hardware.
@@ -1719,17 +1723,19 @@ bitflags::bitflags! {
         /// Read-write depth stencil usage
         const DEPTH_STENCIL_WRITE = 1 << 7;
         /// Read-only storage buffer usage. Corresponds to a UAV in d3d, so is exclusive, despite being read only.
-        const STORAGE_READ = 1 << 8;
-        /// Read-write or write-only storage buffer usage.
+        const STORAGE_READ_ONLY = 1 << 8;
+        /// Write-only storage buffer usage.
+        const STORAGE_WRITE_ONLY = 1 << 9;
+        /// Read-write storage buffer usage.
         const STORAGE_READ_WRITE = 1 << 9;
         /// The combination of states that a texture may be in _at the same time_.
         const INCLUSIVE = Self::COPY_SRC.bits() | Self::RESOURCE.bits() | Self::DEPTH_STENCIL_READ.bits();
         /// The combination of states that a texture must exclusively be in.
-        const EXCLUSIVE = Self::COPY_DST.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ.bits() | Self::STORAGE_READ_WRITE.bits() | Self::PRESENT.bits();
+        const EXCLUSIVE = Self::COPY_DST.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ_ONLY.bits() | Self::STORAGE_READ_WRITE.bits() | Self::PRESENT.bits();
         /// The combination of all usages that the are guaranteed to be be ordered by the hardware.
         /// If a usage is ordered, then if the texture state doesn't change between draw calls, there
         /// are no barriers needed for synchronization.
-        const ORDERED = Self::INCLUSIVE.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ.bits();
+        const ORDERED = Self::INCLUSIVE.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ_ONLY.bits();
 
         /// Flag used by the wgpu-core texture tracker to say a texture is in different states for every sub-resource
         const COMPLEX = 1 << 10;
diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs
index bb7b9926ac..0c00b964a1 100644
--- a/wgpu-hal/src/metal/adapter.rs
+++ b/wgpu-hal/src/metal/adapter.rs
@@ -111,7 +111,9 @@ impl crate::Adapter for super::Adapter {
 
         // Metal defined pixel format capabilities
         let all_caps = Tfc::SAMPLED_LINEAR
-            | Tfc::STORAGE_WRITE
+            | Tfc::STORAGE_READ_ONLY
+            | Tfc::STORAGE_WRITE_ONLY
+            | Tfc::STORAGE_READ_WRITE
             | Tfc::COLOR_ATTACHMENT
             | Tfc::COLOR_ATTACHMENT_BLEND
             | msaa_count
@@ -134,7 +136,7 @@ impl crate::Adapter for super::Adapter {
             | Tf::Rgba8Sint
             | Tf::Rgba16Uint
             | Tf::Rgba16Sint => {
-                read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count
+                read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count
             }
             Tf::R16Unorm
             | Tf::R16Snorm
@@ -143,65 +145,72 @@ impl crate::Adapter for super::Adapter {
             | Tf::Rgba16Unorm
             | Tf::Rgba16Snorm => {
                 Tfc::SAMPLED_LINEAR
-                    | Tfc::STORAGE_WRITE
+                    | Tfc::STORAGE_WRITE_ONLY
                     | Tfc::COLOR_ATTACHMENT
                     | Tfc::COLOR_ATTACHMENT_BLEND
                     | msaa_count
                     | msaa_resolve_desktop_if
             }
             Tf::Rg8Unorm | Tf::Rg16Float | Tf::Bgra8Unorm => all_caps,
-            Tf::Rg8Uint | Tf::Rg8Sint => Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count,
+            Tf::Rg8Uint | Tf::Rg8Sint => {
+                Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count
+            }
             Tf::R32Uint | Tf::R32Sint => {
-                read_write_tier1_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count
+                read_write_tier1_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count
             }
             Tf::R32Float => {
                 let flags = if pc.format_r32float_all {
                     all_caps
                 } else {
-                    Tfc::STORAGE_WRITE
+                    Tfc::STORAGE_WRITE_ONLY
                         | Tfc::COLOR_ATTACHMENT
                         | Tfc::COLOR_ATTACHMENT_BLEND
                         | msaa_count
                 };
                 read_write_tier1_if | flags
             }
-            Tf::Rg16Uint | Tf::Rg16Sint => Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count,
+            Tf::Rg16Uint | Tf::Rg16Sint => {
+                Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count
+            }
             Tf::Rgba8UnormSrgb | Tf::Bgra8UnormSrgb => {
                 let mut flags = all_caps;
-                flags.set(Tfc::STORAGE_WRITE, pc.format_rgba8_srgb_all);
+                flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgba8_srgb_all);
                 flags
             }
             Tf::Rgb10a2Uint => {
                 let mut flags = Tfc::COLOR_ATTACHMENT | msaa_count;
-                flags.set(Tfc::STORAGE_WRITE, pc.format_rgb10a2_uint_write);
+                flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgb10a2_uint_write);
                 flags
             }
             Tf::Rgb10a2Unorm => {
                 let mut flags = all_caps;
-                flags.set(Tfc::STORAGE_WRITE, pc.format_rgb10a2_unorm_all);
+                flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgb10a2_unorm_all);
                 flags
             }
             Tf::Rg11b10Ufloat => {
                 let mut flags = all_caps;
-                flags.set(Tfc::STORAGE_WRITE, pc.format_rg11b10_all);
+                flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rg11b10_all);
                 flags
             }
-            Tf::Rg32Uint | Tf::Rg32Sint => Tfc::COLOR_ATTACHMENT | Tfc::STORAGE_WRITE | msaa_count,
+            Tf::Rg32Uint | Tf::Rg32Sint => {
+                Tfc::COLOR_ATTACHMENT | Tfc::STORAGE_WRITE_ONLY | msaa_count
+            }
             Tf::Rg32Float => {
                 if pc.format_rg32float_all {
                     all_caps
                 } else {
-                    Tfc::STORAGE_WRITE
+                    Tfc::STORAGE_WRITE_ONLY
                         | Tfc::COLOR_ATTACHMENT
                         | Tfc::COLOR_ATTACHMENT_BLEND
                         | msaa_count
                 }
             }
             Tf::Rgba32Uint | Tf::Rgba32Sint => {
-                read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count
+                read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count
             }
             Tf::Rgba32Float => {
-                let mut flags = read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT;
+                let mut flags =
+                    read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT;
                 if pc.format_rgba32float_all {
                     flags |= all_caps
                 } else if pc.msaa_apple7 {
@@ -351,7 +360,7 @@ impl crate::Adapter for super::Adapter {
             usage: crate::TextureUses::COLOR_TARGET
                 | crate::TextureUses::COPY_SRC
                 | crate::TextureUses::COPY_DST
-                | crate::TextureUses::STORAGE_READ
+                //| crate::TextureUses::STORAGE_READ_ONLY
                 | crate::TextureUses::STORAGE_READ_WRITE,
         })
     }
diff --git a/wgpu-hal/src/metal/conv.rs b/wgpu-hal/src/metal/conv.rs
index 6ebabee1a6..f56141d5a7 100644
--- a/wgpu-hal/src/metal/conv.rs
+++ b/wgpu-hal/src/metal/conv.rs
@@ -13,12 +13,12 @@ pub fn map_texture_usage(
     mtl_usage.set(
         metal::MTLTextureUsage::ShaderRead,
         usage.intersects(
-            Tu::RESOURCE | Tu::DEPTH_STENCIL_READ | Tu::STORAGE_READ | Tu::STORAGE_READ_WRITE,
+            Tu::RESOURCE | Tu::DEPTH_STENCIL_READ | Tu::STORAGE_READ_ONLY | Tu::STORAGE_READ_WRITE,
         ),
     );
     mtl_usage.set(
         metal::MTLTextureUsage::ShaderWrite,
-        usage.intersects(Tu::STORAGE_READ_WRITE),
+        usage.intersects(Tu::STORAGE_WRITE_ONLY | Tu::STORAGE_READ_WRITE),
     );
     // needed for combined depth/stencil formats since we might
     // create a stencil-only view from them
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index 40340ca03b..0b6bd1e4a4 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -2115,7 +2115,7 @@ impl crate::Adapter for super::Adapter {
         //     features.contains(vk::FormatFeatureFlags::SAMPLED_IMAGE_FILTER_MINMAX),
         // );
         flags.set(
-            Tfc::STORAGE_WRITE | Tfc::STORAGE_READ_WRITE,
+            Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_WRITE_ONLY | Tfc::STORAGE_READ_ONLY,
             features.contains(vk::FormatFeatureFlags::STORAGE_IMAGE),
         );
         flags.set(
diff --git a/wgpu-hal/src/vulkan/conv.rs b/wgpu-hal/src/vulkan/conv.rs
index 75b807a29c..ce4aa410bf 100644
--- a/wgpu-hal/src/vulkan/conv.rs
+++ b/wgpu-hal/src/vulkan/conv.rs
@@ -263,7 +263,11 @@ pub fn map_texture_usage(usage: crate::TextureUses) -> vk::ImageUsageFlags {
     ) {
         flags |= vk::ImageUsageFlags::DEPTH_STENCIL_ATTACHMENT;
     }
-    if usage.intersects(crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE) {
+    if usage.intersects(
+        crate::TextureUses::STORAGE_READ_ONLY
+            | crate::TextureUses::STORAGE_WRITE_ONLY
+            | crate::TextureUses::STORAGE_READ_WRITE,
+    ) {
         flags |= vk::ImageUsageFlags::STORAGE;
     }
     flags
@@ -305,13 +309,17 @@ pub fn map_texture_usage_to_barrier(
         access |= vk::AccessFlags::DEPTH_STENCIL_ATTACHMENT_READ
             | vk::AccessFlags::DEPTH_STENCIL_ATTACHMENT_WRITE;
     }
-    if usage.contains(crate::TextureUses::STORAGE_READ) {
+    if usage
+        .intersects(crate::TextureUses::STORAGE_READ_ONLY | crate::TextureUses::STORAGE_READ_WRITE)
+    {
         stages |= shader_stages;
         access |= vk::AccessFlags::SHADER_READ;
     }
-    if usage.contains(crate::TextureUses::STORAGE_READ_WRITE) {
+    if usage
+        .intersects(crate::TextureUses::STORAGE_WRITE_ONLY | crate::TextureUses::STORAGE_READ_WRITE)
+    {
         stages |= shader_stages;
-        access |= vk::AccessFlags::SHADER_READ | vk::AccessFlags::SHADER_WRITE;
+        access |= vk::AccessFlags::SHADER_WRITE;
     }
 
     if usage == crate::TextureUses::UNINITIALIZED || usage == crate::TextureUses::PRESENT {
@@ -342,7 +350,7 @@ pub fn map_vk_image_usage(usage: vk::ImageUsageFlags) -> crate::TextureUses {
         bits |= crate::TextureUses::DEPTH_STENCIL_READ | crate::TextureUses::DEPTH_STENCIL_WRITE;
     }
     if usage.contains(vk::ImageUsageFlags::STORAGE) {
-        bits |= crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE;
+        bits |= crate::TextureUses::STORAGE_READ_WRITE;
     }
     bits
 }
@@ -507,7 +515,11 @@ pub fn map_buffer_usage(usage: crate::BufferUses) -> vk::BufferUsageFlags {
     if usage.contains(crate::BufferUses::UNIFORM) {
         flags |= vk::BufferUsageFlags::UNIFORM_BUFFER;
     }
-    if usage.intersects(crate::BufferUses::STORAGE_READ | crate::BufferUses::STORAGE_READ_WRITE) {
+    if usage.intersects(
+        crate::BufferUses::STORAGE_READ_ONLY
+            | crate::BufferUses::STORAGE_WRITE_ONLY
+            | crate::BufferUses::STORAGE_READ_WRITE,
+    ) {
         flags |= vk::BufferUsageFlags::STORAGE_BUFFER;
     }
     if usage.contains(crate::BufferUses::INDEX) {
@@ -561,13 +573,17 @@ pub fn map_buffer_usage_to_barrier(
         stages |= shader_stages;
         access |= vk::AccessFlags::UNIFORM_READ;
     }
-    if usage.intersects(crate::BufferUses::STORAGE_READ) {
+    if usage
+        .intersects(crate::BufferUses::STORAGE_READ_ONLY | crate::BufferUses::STORAGE_READ_WRITE)
+    {
         stages |= shader_stages;
         access |= vk::AccessFlags::SHADER_READ;
     }
-    if usage.intersects(crate::BufferUses::STORAGE_READ_WRITE) {
+    if usage
+        .intersects(crate::BufferUses::STORAGE_WRITE_ONLY | crate::BufferUses::STORAGE_READ_WRITE)
+    {
         stages |= shader_stages;
-        access |= vk::AccessFlags::SHADER_READ | vk::AccessFlags::SHADER_WRITE;
+        access |= vk::AccessFlags::SHADER_WRITE;
     }
     if usage.contains(crate::BufferUses::INDEX) {
         stages |= vk::PipelineStageFlags::VERTEX_INPUT;
diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs
index 546a3fde0d..43676a6822 100644
--- a/wgpu-types/src/lib.rs
+++ b/wgpu-types/src/lib.rs
@@ -2381,8 +2381,14 @@ bitflags::bitflags! {
         /// to a render pass for an automatic driver-implemented resolve.
         const MULTISAMPLE_RESOLVE = 1 << 5;
         /// When used as a STORAGE texture, then a texture with this format can be bound with
-        /// [`StorageTextureAccess::ReadOnly`] or [`StorageTextureAccess::ReadWrite`].
-        const STORAGE_WRITE = 1 << 6;
+        /// [`StorageTextureAccess::ReadOnly`].
+        const STORAGE_READ_ONLY = 1 << 8;
+        /// When used as a STORAGE texture, then a texture with this format can be bound with
+        /// [`StorageTextureAccess::WriteOnly`].
+        const STORAGE_WRITE_ONLY = 1 << 6;
+        /// When used as a STORAGE texture, then a texture with this format can be bound with
+        /// any [`StorageTextureAccess`].
+        const STORAGE_READ_WRITE = 1 << 9;
         /// If not present, the texture can't be blended into the render target.
         const BLENDABLE = 1 << 7;
     }