diff --git a/src/dawn_native/QueryHelper.cpp b/src/dawn_native/QueryHelper.cpp
index 70d9a9c..eb3d0bf 100644
--- a/src/dawn_native/QueryHelper.cpp
+++ b/src/dawn_native/QueryHelper.cpp
@@ -55,9 +55,9 @@
             };
 
             [[group(0), binding(0)]]
-                var<storage> timestamps : [[access(read_write)]] TimestampArr;
+                var<storage, read_write> timestamps : TimestampArr;
             [[group(0), binding(1)]]
-                var<storage> availability : [[access(read)]] AvailabilityArr;
+                var<storage, read> availability : AvailabilityArr;
             [[group(0), binding(2)]] var<uniform> params : TimestampParams;
 
 
diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
index ed4727d..f7a5a91 100644
--- a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
@@ -234,7 +234,7 @@
                      bgl->GetBindingInfo(bindingIndex).buffer.type ==
                          wgpu::BufferBindingType::Storage);
                 if (forceStorageBufferAsUAV) {
-                    accessControls.emplace(srcBindingPoint, tint::ast::AccessControl::kReadWrite);
+                    accessControls.emplace(srcBindingPoint, tint::ast::Access::kReadWrite);
                 }
             }
         }
diff --git a/src/tests/DawnTest.cpp b/src/tests/DawnTest.cpp
index 0d972fc..3de7da0 100644
--- a/src/tests/DawnTest.cpp
+++ b/src/tests/DawnTest.cpp
@@ -1098,7 +1098,7 @@
         };
 
         [[group(0), binding(0)]] var tex : texture_depth_2d;
-        [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
+        [[group(0), binding(1)]] var<storage, read_write> result : Result;
 
         [[stage(compute)]] fn main(
             [[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>
diff --git a/src/tests/end2end/BindGroupTests.cpp b/src/tests/end2end/BindGroupTests.cpp
index fbe6e52..4ab61c4 100644
--- a/src/tests/end2end/BindGroupTests.cpp
+++ b/src/tests/end2end/BindGroupTests.cpp
@@ -72,8 +72,8 @@
                        << " : Buffer" << i << ";";
                     break;
                 case wgpu::BufferBindingType::Storage:
-                    fs << "\n[[group(" << i << "), binding(0)]] var<storage> buffer" << i
-                       << " : [[access(read)]] Buffer" << i << ";";
+                    fs << "\n[[group(" << i << "), binding(0)]] var<storage, read> buffer" << i
+                       << " : Buffer" << i << ";";
                     break;
                 default:
                     UNREACHABLE();
@@ -837,9 +837,9 @@
         };
 
         [[group(0), binding(2)]] var<uniform> buffer2 : Buffer2;
-        [[group(0), binding(3)]] var<storage> buffer3 : [[access(read)]] Buffer3;
-        [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0;
-        [[group(0), binding(4)]] var<storage> outputBuffer : [[access(read_write)]] OutputBuffer;
+        [[group(0), binding(3)]] var<storage, read> buffer3 : Buffer3;
+        [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0;
+        [[group(0), binding(4)]] var<storage, read_write> outputBuffer : OutputBuffer;
 
         [[stage(compute)]] fn main() {
             outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value);
@@ -1103,7 +1103,7 @@
         [[block]] struct Buffer0 {
             color : vec4<f32>;
         };
-        [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0;
+        [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0;
 
         [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
             return buffer0.color;
@@ -1240,8 +1240,7 @@
             };
         )";
         interface << "[[group(0), binding(" << binding++ << ")]] "
-                  << "var<storage> sbuf" << i << " : [[access(read)]] ReadOnlyStorageBuffer" << i
-                  << ";\n";
+                  << "var<storage, read> sbuf" << i << " : ReadOnlyStorageBuffer" << i << ";\n";
 
         body << "if (sbuf" << i << ".value != " << expectedValue++ << "u) {\n";
         body << "    return;\n";
@@ -1257,7 +1256,7 @@
         };
     )";
     interface << "[[group(0), binding(" << binding++ << ")]] "
-              << "var<storage> result : [[access(read_write)]] ReadWriteStorageBuffer;\n";
+              << "var<storage, read_write> result : ReadWriteStorageBuffer;\n";
 
     body << "result.value = 1u;\n";
 
diff --git a/src/tests/end2end/BufferZeroInitTests.cpp b/src/tests/end2end/BufferZeroInitTests.cpp
index 55bdf90..78e1478 100644
--- a/src/tests/end2end/BufferZeroInitTests.cpp
+++ b/src/tests/end2end/BufferZeroInitTests.cpp
@@ -1035,7 +1035,7 @@
         [[block]] struct SSBO {
             value : vec4<u32>;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
         [[group(0), binding(1)]] var outImage : texture_storage_2d<rgba8unorm, write>;
 
         [[stage(compute)]] fn main() {
@@ -1074,7 +1074,7 @@
         [[block]] struct SSBO {
             value : array<vec4<u32>, 2>;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
         [[group(0), binding(1)]] var outImage : texture_storage_2d<rgba8unorm, write>;
 
         [[stage(compute)]] fn main() {
diff --git a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
index f41da7b..721ff21 100644
--- a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
+++ b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
@@ -96,8 +96,8 @@
         };
 
         // TODO(crbug.com/tint/386): Use the same struct type
-        [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
-        [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
+        [[set(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -123,8 +123,8 @@
         };
 
         // TODO(crbug.com/tint/386): Use the same struct type
-        [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
-        [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
+        [[set(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -145,8 +145,8 @@
         };
 
         // TODO(crbug.com/tint/386): Use the same struct type
-        [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
-        [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
+        [[set(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
diff --git a/src/tests/end2end/ComputeDispatchTests.cpp b/src/tests/end2end/ComputeDispatchTests.cpp
index 93d844f..860aaee 100644
--- a/src/tests/end2end/ComputeDispatchTests.cpp
+++ b/src/tests/end2end/ComputeDispatchTests.cpp
@@ -37,7 +37,7 @@
             };
 
             [[group(0), binding(0)]] var<uniform> input : InputBuf;
-            [[group(0), binding(1)]] var<storage> output : [[access(read_write)]] OutputBuf;
+            [[group(0), binding(1)]] var<storage, read_write> output : OutputBuf;
 
             [[stage(compute), workgroup_size(1, 1, 1)]]
             fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
diff --git a/src/tests/end2end/ComputeSharedMemoryTests.cpp b/src/tests/end2end/ComputeSharedMemoryTests.cpp
index 8ac3212..3e7e6de 100644
--- a/src/tests/end2end/ComputeSharedMemoryTests.cpp
+++ b/src/tests/end2end/ComputeSharedMemoryTests.cpp
@@ -78,7 +78,7 @@
             x : u32;
         };
 
-        [[group(0), binding(0)]] var<storage> dst : [[access(write)]] Dst;
+        [[group(0), binding(0)]] var<storage, write> dst : Dst;
         var<workgroup> tmp : u32;
 
         [[stage(compute), workgroup_size(4,4,1)]]
diff --git a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
index 3cfa6dc..7775974 100644
--- a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
+++ b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
@@ -36,7 +36,7 @@
             data : array<u32, 100>;
         };
 
-        [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf;
+        [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -91,8 +91,8 @@
             data : array<u32, 100>;
         };
 
-        [[group(0), binding(0)]] var<storage> src : [[access(read_write)]] Src;
-        [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst;
+        [[group(0), binding(0)]] var<storage, read_write> src : Src;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Dst;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -162,8 +162,8 @@
             data : array<u32, 100>;
         };
 
-        [[group(0), binding(0)]] var<storage> src : [[access(read)]] Src;
-        [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst;
+        [[group(0), binding(0)]] var<storage, read> src : Src;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Dst;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -231,7 +231,7 @@
         };
 
         [[group(0), binding(0)]] var<uniform> src : Buf;
-        [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -299,7 +299,7 @@
         };
 
         [[group(0), binding(0)]] var<uniform> src : Buf;
-        [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
 
         [[stage(compute)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -359,7 +359,7 @@
         [[block]] struct Buf {
             data : array<u32, 3>;
         };
-        [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf;
+        [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
 
         [[stage(compute)]] fn main() {
             buf.data = array<u32, 3>(1u, 1u, 1u);
@@ -373,12 +373,12 @@
         [[block]] struct Buf {
             data : array<u32, 3>;
         };
-        [[group(0), binding(0)]] var<storage> buf : [[access(read)]] Buf;
+        [[group(0), binding(0)]] var<storage, read> buf : Buf;
 
         [[block]] struct Result {
             data : u32;
         };
-        [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
+        [[group(0), binding(1)]] var<storage, read_write> result : Result;
 
         [[stage(compute)]] fn main() {
             result.data = 2u;
diff --git a/src/tests/end2end/CopyTextureForBrowserTests.cpp b/src/tests/end2end/CopyTextureForBrowserTests.cpp
index ffcf2fc..1cce29a 100644
--- a/src/tests/end2end/CopyTextureForBrowserTests.cpp
+++ b/src/tests/end2end/CopyTextureForBrowserTests.cpp
@@ -146,7 +146,7 @@
             };
             [[group(0), binding(0)]] var src : texture_2d<f32>;
             [[group(0), binding(1)]] var dst : texture_2d<f32>;
-            [[group(0), binding(2)]] var<storage> output : [[access(read_write)]] OutputBuf;
+            [[group(0), binding(2)]] var<storage, read_write> output : OutputBuf;
             [[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
             fn aboutEqual(value : f32, expect : f32) -> bool {
                 // The value diff should be smaller than the hard coded tolerance.
@@ -389,7 +389,6 @@
 // Verify CopyTextureForBrowserTests works with internal pipeline.
 // The case do copy without any transform.
 TEST_P(CopyTextureForBrowserTests, PassthroughCopy) {
-
     constexpr uint32_t kWidth = 10;
     constexpr uint32_t kHeight = 1;
 
@@ -420,7 +419,6 @@
 }
 
 TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) {
-
     constexpr uint32_t kWidth = 899;
     constexpr uint32_t kHeight = 999;
 
diff --git a/src/tests/end2end/CreatePipelineAsyncTests.cpp b/src/tests/end2end/CreatePipelineAsyncTests.cpp
index 44990e0..c5e76f8 100644
--- a/src/tests/end2end/CreatePipelineAsyncTests.cpp
+++ b/src/tests/end2end/CreatePipelineAsyncTests.cpp
@@ -78,7 +78,7 @@
         [[block]] struct SSBO {
             value : u32;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
 
         [[stage(compute)]] fn main() {
             ssbo.value = 1u;
@@ -113,7 +113,7 @@
         [[block]] struct SSBO {
             value : u32;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
 
         [[stage(compute)]] fn main() {
             ssbo.value = 1u;
@@ -312,7 +312,7 @@
         [[block]] struct SSBO {
             value : u32;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
 
         [[stage(compute)]] fn main() {
             ssbo.value = 1u;
@@ -353,7 +353,7 @@
         [[block]] struct SSBO {
             value : u32;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
 
         [[stage(compute)]] fn main() {
             ssbo.value = 1u;
diff --git a/src/tests/end2end/D3D12CachingTests.cpp b/src/tests/end2end/D3D12CachingTests.cpp
index fde7d6c..ecf7c8b 100644
--- a/src/tests/end2end/D3D12CachingTests.cpp
+++ b/src/tests/end2end/D3D12CachingTests.cpp
@@ -211,7 +211,7 @@
         [[block]] struct Data {
             data : u32;
         };
-        [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data;
+        [[binding(0), group(0)]] var<storage, read_write> data : Data;
 
         [[stage(compute)]] fn write1() {
             data.data = 1u;
diff --git a/src/tests/end2end/DepthStencilSamplingTests.cpp b/src/tests/end2end/DepthStencilSamplingTests.cpp
index 1b7e584..d0681a0 100644
--- a/src/tests/end2end/DepthStencilSamplingTests.cpp
+++ b/src/tests/end2end/DepthStencilSamplingTests.cpp
@@ -141,8 +141,8 @@
                                  << " : texture_2d<f32>;\n";
 
                     shaderSource << "[[group(0), binding(" << 2 * index + 1
-                                 << ")]] var<storage> result" << index
-                                 << " : [[access(read_write)]] DepthResult;\n";
+                                 << ")]] var<storage, read_write> result" << index
+                                 << " : DepthResult;\n";
 
                     shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index
                                << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];";
@@ -152,8 +152,8 @@
                                  << " : texture_2d<u32>;\n";
 
                     shaderSource << "[[group(0), binding(" << 2 * index + 1
-                                 << ")]] var<storage> result" << index
-                                 << " : [[access(read_write)]] StencilResult;\n";
+                                 << ")]] var<storage, read_write> result" << index
+                                 << " : StencilResult;\n";
 
                     shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index
                                << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];";
@@ -221,7 +221,7 @@
             [[block]] struct SamplerResult {
                 value : f32;
             };
-            [[group(0), binding(3)]] var<storage> samplerResult : [[access(read_write)]] SamplerResult;
+            [[group(0), binding(3)]] var<storage, read_write> samplerResult : SamplerResult;
 
             [[stage(compute)]] fn main() {
                 samplerResult.value = textureSampleCompare(tex, samp, vec2<f32>(0.5, 0.5), uniforms.compareRef);
diff --git a/src/tests/end2end/DynamicBufferOffsetTests.cpp b/src/tests/end2end/DynamicBufferOffsetTests.cpp
index 33de9ef..32821fd 100644
--- a/src/tests/end2end/DynamicBufferOffsetTests.cpp
+++ b/src/tests/end2end/DynamicBufferOffsetTests.cpp
@@ -125,9 +125,9 @@
             };
 
             [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
-            [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2;
+            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2;
             [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
-            [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4;
+            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4;
         )";
 
         if (isInheritedPipeline) {
@@ -192,9 +192,9 @@
             };
 
             [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
-            [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2;
+            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2;
             [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
-            [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4;
+            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4;
         )";
 
         if (isInheritedPipeline) {
diff --git a/src/tests/end2end/EntryPointTests.cpp b/src/tests/end2end/EntryPointTests.cpp
index 037d6ed..da5d5db 100644
--- a/src/tests/end2end/EntryPointTests.cpp
+++ b/src/tests/end2end/EntryPointTests.cpp
@@ -64,7 +64,7 @@
         [[block]] struct Data {
             data : u32;
         };
-        [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data;
+        [[binding(0), group(0)]] var<storage, read_write> data : Data;
 
         [[stage(compute)]] fn write1() {
             data.data = 1u;
diff --git a/src/tests/end2end/FirstIndexOffsetTests.cpp b/src/tests/end2end/FirstIndexOffsetTests.cpp
index b6e56d7..506c0d3 100644
--- a/src/tests/end2end/FirstIndexOffsetTests.cpp
+++ b/src/tests/end2end/FirstIndexOffsetTests.cpp
@@ -132,7 +132,7 @@
   vertex_index : u32;
   instance_index : u32;
 };
-[[group(0), binding(0)]] var<storage> idx_vals : [[access(read_write)]] IndexVals;
+[[group(0), binding(0)]] var<storage, read_write> idx_vals : IndexVals;
 
 struct FragInputs {
 )" + fragmentInputs.str() + R"(
diff --git a/src/tests/end2end/GpuMemorySynchronizationTests.cpp b/src/tests/end2end/GpuMemorySynchronizationTests.cpp
index 1db5400..7113a13 100644
--- a/src/tests/end2end/GpuMemorySynchronizationTests.cpp
+++ b/src/tests/end2end/GpuMemorySynchronizationTests.cpp
@@ -39,7 +39,7 @@
             [[block]] struct Data {
                 a : i32;
             };
-            [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
+            [[group(0), binding(0)]] var<storage, read_write> data : Data;
             [[stage(compute)]] fn main() {
                 data.a = data.a + 1;
             })");
@@ -66,7 +66,7 @@
             [[block]] struct Data {
                 i : i32;
             };
-            [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
+            [[group(0), binding(0)]] var<storage, read_write> data : Data;
             [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
                 data.i = data.i + 1;
                 return vec4<f32>(f32(data.i) / 255.0, 0.0, 0.0, 1.0);
@@ -255,7 +255,7 @@
             sampledOut: u32;
             storageOut: u32;
         };
-        [[group(0), binding(0)]] var<storage> output : [[access(write)]] Output;
+        [[group(0), binding(0)]] var<storage, write> output : Output;
         [[group(0), binding(1)]] var sampledTex : texture_2d<u32>;
         [[group(0), binding(2)]] var storageTex : texture_storage_2d<r32uint, read>;
 
@@ -315,7 +315,7 @@
             [[block]] struct Data {
                 a : f32;
             };
-            [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
+            [[group(0), binding(0)]] var<storage, read_write> data : Data;
             [[stage(compute)]] fn main() {
                 data.a = 1.0;
             })");
@@ -512,12 +512,12 @@
         [[block]] struct VBContents {
             pos : array<vec4<f32>, 4>;
         };
-        [[group(0), binding(0)]] var<storage> vbContents : [[access(read_write)]] VBContents;
+        [[group(0), binding(0)]] var<storage, read_write> vbContents : VBContents;
 
         [[block]] struct IBContents {
             indices : array<vec4<i32>, 2>;
         };
-        [[group(0), binding(1)]] var<storage> ibContents : [[access(read_write)]] IBContents;
+        [[group(0), binding(1)]] var<storage, read_write> ibContents : IBContents;
 
         // TODO(crbug.com/tint/386): Use the same struct.
         [[block]] struct ColorContents1 {
@@ -526,8 +526,8 @@
         [[block]] struct ColorContents2 {
             color : f32;
         };
-        [[group(0), binding(2)]] var<storage> uniformContents : [[access(read_write)]] ColorContents1;
-        [[group(0), binding(3)]] var<storage> storageContents : [[access(read_write)]] ColorContents2;
+        [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents1;
+        [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents2;
 
         [[stage(compute)]] fn main() {
             vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
@@ -581,7 +581,7 @@
         };
 
         [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
-        [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf;
+        [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf;
 
         [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
             return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0);
@@ -642,7 +642,7 @@
             [[align(256)]] color1 : f32;
         };
 
-        [[group(0), binding(0)]] var<storage> contents : [[access(read_write)]] Contents;
+        [[group(0), binding(0)]] var<storage, read_write> contents : Contents;
 
         [[stage(compute)]] fn main() {
             contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
@@ -696,7 +696,7 @@
             color : f32;
         };
         [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
-        [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf;
+        [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf;
 
         [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
             return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0);
diff --git a/src/tests/end2end/MultisampledSamplingTests.cpp b/src/tests/end2end/MultisampledSamplingTests.cpp
index 055f2c1..39e6103 100644
--- a/src/tests/end2end/MultisampledSamplingTests.cpp
+++ b/src/tests/end2end/MultisampledSamplingTests.cpp
@@ -100,7 +100,7 @@
                     colorSamples : array<f32, 4>;
                     depthSamples : array<f32, 4>;
                 };
-                [[group(0), binding(2)]] var<storage> results : [[access(read_write)]] Results;
+                [[group(0), binding(2)]] var<storage, read_write> results : Results;
 
                 [[stage(compute)]] fn main() {
                     for (var i : i32 = 0; i < 4; i = i + 1) {
diff --git a/src/tests/end2end/OpArrayLengthTests.cpp b/src/tests/end2end/OpArrayLengthTests.cpp
index c9d42f7..b7b55bf 100644
--- a/src/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/tests/end2end/OpArrayLengthTests.cpp
@@ -63,10 +63,10 @@
             };
 
             // The length should be 1 because the buffer is 4-byte long.
-            [[group(0), binding(0)]] var<storage> buffer1 : [[access(read)]] DataBuffer1;
+            [[group(0), binding(0)]] var<storage, read> buffer1 : DataBuffer1;
 
             // The length should be 64 because the buffer is 256 bytes long.
-            [[group(0), binding(1)]] var<storage> buffer2 : [[access(read)]] DataBuffer2;
+            [[group(0), binding(1)]] var<storage, read> buffer2 : DataBuffer2;
 
             // The length should be (512 - 16*4) / 8 = 56 because the buffer is 512 bytes long
             // and the structure is 8 bytes big.
@@ -79,7 +79,7 @@
                 [[size(64)]] garbage : mat4x4<f32>;
                 data : [[stride(8)]] array<Buffer3Data>;
             };
-            [[group(0), binding(2)]] var<storage> buffer3 : [[access(read)]] Buffer3;
+            [[group(0), binding(2)]] var<storage, read> buffer3 : Buffer3;
         )";
 
         // See comments in the shader for an explanation of these values
@@ -128,7 +128,7 @@
         [[block]] struct ResultBuffer {
             data : [[stride(4)]] array<u32, 3>;
         };
-        [[group(1), binding(0)]] var<storage> result : [[access(read_write)]] ResultBuffer;
+        [[group(1), binding(0)]] var<storage, read_write> result : ResultBuffer;
         )" + mShaderInterface + R"(
         [[stage(compute)]] fn main() {
             result.data[0] = arrayLength(buffer1.data);
diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp
index 7c9475c..ace5c8f 100644
--- a/src/tests/end2end/ShaderTests.cpp
+++ b/src/tests/end2end/ShaderTests.cpp
@@ -35,7 +35,7 @@
     data : array<u32, 19>;
 };
 
-[[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf;
+[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
 
 [[stage(compute)]] fn main() {
     let factor : f32 = 1.0001;
diff --git a/src/tests/end2end/StorageTextureTests.cpp b/src/tests/end2end/StorageTextureTests.cpp
index 0ab4a0a..a25aa53 100644
--- a/src/tests/end2end/StorageTextureTests.cpp
+++ b/src/tests/end2end/StorageTextureTests.cpp
@@ -167,12 +167,12 @@
         std::ostringstream ostream;
         ostream << "[[group(0), binding(" << binding << ")]] "
                 << "var storageImage" << binding << " : "
-                << "[[access(" << accessQualifier << ")]] "
                 << "texture_storage_2d";
         if (is2DArray) {
             ostream << "_array";
         }
-        ostream << "<" << utils::GetWGSLImageFormatQualifier(format) << ">;";
+        ostream << "<" << utils::GetWGSLImageFormatQualifier(format) << ", ";
+        ostream << accessQualifier << ">;";
         return ostream.str();
     }
 
@@ -710,7 +710,7 @@
   result : u32;
 };
 
-[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
+[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer;
 )" << CommonReadOnlyTestCode(format)
                  << R"(
 [[stage(compute)]] fn main() {
@@ -934,7 +934,7 @@
   result : u32;
 };
 
-[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
+[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer;
 )" << CommonReadOnlyTestCode(kTextureFormat, true)
              << R"(
 [[stage(compute)]] fn main() {
@@ -1202,7 +1202,7 @@
 };
 
 [[group(0), binding(0)]] var srcImage : texture_storage_2d<r32uint, read>;
-[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
+[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer;
 )") + kCommonReadOnlyZeroInitTestCode + R"(
 [[stage(compute)]] fn main() {
   if (doTest()) {
diff --git a/src/tests/end2end/TextureZeroInitTests.cpp b/src/tests/end2end/TextureZeroInitTests.cpp
index 3f2e8ee..f8a03d5 100644
--- a/src/tests/end2end/TextureZeroInitTests.cpp
+++ b/src/tests/end2end/TextureZeroInitTests.cpp
@@ -979,7 +979,7 @@
         [[block]] struct Result {
             value : vec4<f32>;
         };
-        [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
+        [[group(0), binding(1)]] var<storage, read_write> result : Result;
         [[stage(compute)]] fn main() {
            result.value = textureLoad(tex, vec2<i32>(0,0), 0);
         }
diff --git a/src/tests/perf_tests/ShaderRobustnessPerf.cpp b/src/tests/perf_tests/ShaderRobustnessPerf.cpp
index 0add757..f63af7b 100644
--- a/src/tests/perf_tests/ShaderRobustnessPerf.cpp
+++ b/src/tests/perf_tests/ShaderRobustnessPerf.cpp
@@ -29,9 +29,9 @@
             numbers: array<f32>;
         };
 
-        [[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix;
-        [[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix;
-        [[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix;
+        [[group(0), binding(0)]] var<storage, read> firstMatrix : Matrix;
+        [[group(0), binding(1)]] var<storage, read> secondMatrix : Matrix;
+        [[group(0), binding(2)]] var<storage, write> resultMatrix : Matrix;
         [[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
 
         fn mm_readA(row : u32, col : u32) -> f32  {
@@ -196,9 +196,9 @@
             numbers: array<vec4<f32>>;
         };
 
-        [[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix;
-        [[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix;
-        [[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix;
+        [[group(0), binding(0)]] var<storage, read> firstMatrix : Matrix;
+        [[group(0), binding(1)]] var<storage, read> secondMatrix : Matrix;
+        [[group(0), binding(2)]] var<storage, write> resultMatrix : Matrix;
         [[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
 
         fn mm_readA(row : u32, col : u32) -> vec4<f32>  {
diff --git a/src/tests/unittests/validation/BindGroupValidationTests.cpp b/src/tests/unittests/validation/BindGroupValidationTests.cpp
index 33c656f..b41b8c4 100644
--- a/src/tests/unittests/validation/BindGroupValidationTests.cpp
+++ b/src/tests/unittests/validation/BindGroupValidationTests.cpp
@@ -1278,8 +1278,8 @@
 
                 [[group(0), binding(0)]] var<uniform> uBufferDynamic : S;
                 [[group(0), binding(1)]] var<uniform> uBuffer : S;
-                [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
-                [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
+                [[group(0), binding(2)]] var<storage, read_write> sBufferDynamic : S;
+                [[group(0), binding(3)]] var<storage, read> sReadonlyBufferDynamic : S;
 
                 [[stage(fragment)]] fn main() {
                 })");
@@ -1301,8 +1301,8 @@
 
                 [[group(0), binding(0)]] var<uniform> uBufferDynamic : S;
                 [[group(0), binding(1)]] var<uniform> uBuffer : S;
-                [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
-                [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
+                [[group(0), binding(2)]] var<storage, read_write> sBufferDynamic : S;
+                [[group(0), binding(3)]] var<storage, read> sReadonlyBufferDynamic : S;
 
                 [[stage(compute), workgroup_size(4, 4, 1)]] fn main() {
                 })");
@@ -1728,8 +1728,7 @@
                 ss << "[[group(" << l << "), binding(" << b << ")]] ";
                 switch (binding) {
                     case wgpu::BufferBindingType::Storage:
-                        ss << "var<storage> set" << l << "_binding" << b
-                           << " : [[access(read_write)]] S;";
+                        ss << "var<storage, read_write> set" << l << "_binding" << b << " : S;";
                         break;
                     case wgpu::BufferBindingType::Uniform:
                         ss << "var<uniform> set" << l << "_binding" << b << " : S;";
@@ -1897,8 +1896,8 @@
                 value : vec2<f32>;
             };
 
-            [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
-            [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
+            [[group(0), binding(0)]] var<storage, read_write> sBufferDynamic : S;
+            [[group(1), binding(0)]] var<storage, read> sReadonlyBufferDynamic : S;
 
             [[stage(fragment)]] fn main() {
                 var val : vec2<f32> = sBufferDynamic.value;
@@ -1932,8 +1931,8 @@
                 value : vec2<f32>;
             };
 
-            [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
-            [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
+            [[group(0), binding(0)]] var<storage, read_write> sBufferDynamic : S;
+            [[group(1), binding(0)]] var<storage, read> sReadonlyBufferDynamic : S;
 
             [[stage(compute), workgroup_size(4, 4, 1)]] fn main() {
                 var val : vec2<f32> = sBufferDynamic.value;
diff --git a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
index de024b0..a0a531a 100644
--- a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
+++ b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
@@ -66,7 +66,7 @@
         [[block]] struct S3 {
             pos : mat4x4<f32>;
         };
-        [[group(3), binding(0)]] var<storage> storage3 : [[access(read_write)]] S3;
+        [[group(3), binding(0)]] var<storage, read_write> storage3 : S3;
 
         [[stage(fragment)]] fn main() {
             var pos_u : vec4<f32> = uniform2.pos;
@@ -202,7 +202,7 @@
             [[block]] struct S {
                 pos : vec4<f32>;
             };
-            [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S;
+            [[group(0), binding(0)]] var<storage, read_write> ssbo : S;
 
             [[stage(fragment)]] fn main() {
                 var pos : vec4<f32> = ssbo.pos;
@@ -229,7 +229,7 @@
             [[block]] struct S {
                 pos : vec4<f32>;
             };
-            [[group(0), binding(0)]] var<storage> ssbo : [[access(read)]] S;
+            [[group(0), binding(0)]] var<storage, read> ssbo : S;
 
             [[stage(fragment)]] fn main() {
                 var pos : vec4<f32> = ssbo.pos;
@@ -707,7 +707,7 @@
         [[block]] struct S {
             pos : vec4<f32>;
         };
-        [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S;
+        [[group(0), binding(0)]] var<storage, read_write> ssbo : S;
 
         [[stage(fragment)]] fn main() {
             var pos : vec4<f32> = ssbo.pos;
@@ -912,8 +912,8 @@
         [[block]] struct Data {
             data : f32;
         };
-        [[group(0), binding(0)]] var<storage> data0 : [[access(read_write)]] Data;
-        [[group(0), binding(1)]] var<storage> data1 : [[access(read_write)]] Data;
+        [[group(0), binding(0)]] var<storage, read_write> data0 : Data;
+        [[group(0), binding(1)]] var<storage, read_write> data1 : Data;
 
         [[stage(compute)]] fn compute0() {
             data0.data = 0.0;
diff --git a/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp b/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
index de4ec17..df88025 100644
--- a/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
+++ b/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
@@ -80,12 +80,10 @@
                     ostream << "var<uniform> b" << index << " : S" << index << ";\n";
                     break;
                 case wgpu::BufferBindingType::Storage:
-                    ostream << "var<storage> b" << index << " : [[access(read_write)]] S" << index
-                            << ";\n";
+                    ostream << "var<storage, read_write> b" << index << " : S" << index << ";\n";
                     break;
                 case wgpu::BufferBindingType::ReadOnlyStorage:
-                    ostream << "var<storage> b" << index << " : [[access(read)]] S" << index
-                            << ";\n";
+                    ostream << "var<storage, read> b" << index << " : S" << index << ";\n";
                     break;
                 default:
                     UNREACHABLE();
diff --git a/src/tests/unittests/validation/RenderBundleValidationTests.cpp b/src/tests/unittests/validation/RenderBundleValidationTests.cpp
index 11277b2..b6da94e 100644
--- a/src/tests/unittests/validation/RenderBundleValidationTests.cpp
+++ b/src/tests/unittests/validation/RenderBundleValidationTests.cpp
@@ -46,7 +46,7 @@
                 [[block]] struct Storage {
                     dummy : array<f32>;
                 };
-                [[group(1), binding(1)]] var<storage> ssbo : [[access(read_write)]] Storage;
+                [[group(1), binding(1)]] var<storage, read_write> ssbo : Storage;
 
                 [[stage(fragment)]] fn main() {
                 })");
diff --git a/src/tests/unittests/validation/RenderPipelineValidationTests.cpp b/src/tests/unittests/validation/RenderPipelineValidationTests.cpp
index 37f156f..593a365 100644
--- a/src/tests/unittests/validation/RenderPipelineValidationTests.cpp
+++ b/src/tests/unittests/validation/RenderPipelineValidationTests.cpp
@@ -460,7 +460,7 @@
         [[block]] struct Dst {
             data : array<u32, 100>;
         };
-        [[group(0), binding(0)]] var<storage> dst : [[access(read_write)]] Dst;
+        [[group(0), binding(0)]] var<storage, read_write> dst : Dst;
         [[stage(vertex)]] fn main([[builtin(vertex_index)]] VertexIndex : u32) -> [[builtin(position)]] vec4<f32> {
             dst.data[VertexIndex] = 0x1234u;
             return vec4<f32>();
diff --git a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
index 142ba56..d70ef9a 100644
--- a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
+++ b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
@@ -762,7 +762,7 @@
                 [[block]] struct RBuffer {
                     value : f32;
                 };
-                [[group(0), binding(0)]] var<storage> rBuffer : [[access(read)]] RBuffer;
+                [[group(0), binding(0)]] var<storage, read> rBuffer : RBuffer;
                 [[stage(fragment)]] fn main() {
                 })");
             utils::ComboRenderPipelineDescriptor pipelineDescriptor;
diff --git a/src/tests/unittests/validation/StorageTextureValidationTests.cpp b/src/tests/unittests/validation/StorageTextureValidationTests.cpp
index 5294c0f..a212e10 100644
--- a/src/tests/unittests/validation/StorageTextureValidationTests.cpp
+++ b/src/tests/unittests/validation/StorageTextureValidationTests.cpp
@@ -82,9 +82,8 @@
         }
 
         std::ostringstream ostream;
-        ostream << "[[group(0), binding(0)]] var image0 : "
-                << "[[access(" << access << ")]] " << imageTypeDeclaration << "<"
-                << imageFormatQualifier
+        ostream << "[[group(0), binding(0)]] var image0 : " << imageTypeDeclaration << "<"
+                << imageFormatQualifier << ", " << access
                 << ">;\n"
                    "[[stage(compute)]] fn main() {\n"
                    "    textureDimensions(image0);\n"
@@ -193,7 +192,7 @@
             [[block]] struct Buf {
                 data : f32;
             };
-            [[group(0), binding(1)]] var<storage> buf : [[access(read_write)]] Buf;
+            [[group(0), binding(1)]] var<storage, read_write> buf : Buf;
 
             [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
                  buf.data = textureLoad(image0, vec2<i32>(LocalInvocationID.xy)).x;