tint: Disallow write-only storage buffers

These have not been in the spec for a long time. The read_write access
mode can be used instead.

Fixed: tint:1342
Change-Id: I01ffc343d2d2f9df9d7028bba4548c749616c65c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/93500
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>

src/dawn/native/ComputePassEncoder.cpp

@@ -62,7 +62,7 @@ ResultOrError<ComputePipelineBase*> GetOrCreateIndirectDispatchValidationPipelin
 
                 @group(0) @binding(0) var<uniform> uniformParams: UniformParams;
                 @group(0) @binding(1) var<storage, read_write> clientParams: IndirectParams;
-                @group(0) @binding(2) var<storage, write> validatedParams: ValidatedParams;
+                @group(0) @binding(2) var<storage, read_write> validatedParams: ValidatedParams;
 
                 @compute @workgroup_size(1, 1, 1)
                 fn main() {

src/dawn/native/IndirectDrawValidationEncoder.cpp

@@ -81,7 +81,7 @@ static const char sRenderValidationShaderSource[] = R"(
 
             @group(0) @binding(0) var<storage, read> batch: BatchInfo;
             @group(0) @binding(1) var<storage, read_write> inputParams: IndirectParams;
-            @group(0) @binding(2) var<storage, write> outputParams: IndirectParams;
+            @group(0) @binding(2) var<storage, read_write> outputParams: IndirectParams;
 
             fn numIndirectParamsPerDrawCallInput() -> u32 {
                 var numParams = kNumDrawIndirectParams;

src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp

@@ -78,7 +78,7 @@ TEST_P(ComputeSharedMemoryTests, Basic) {
             x : u32
         }
 
-        @group(0) @binding(0) var<storage, write> dst : Dst;
+        @group(0) @binding(0) var<storage, read_write> dst : Dst;
         var<workgroup> tmp : u32;
 
         @compute @workgroup_size(4,4,1)
@@ -117,7 +117,7 @@ TEST_P(ComputeSharedMemoryTests, AssortedTypes) {
             d_vector : vec4<f32>,
         }
 
-        @group(0) @binding(0) var<storage, write> dst : Dst;
+        @group(0) @binding(0) var<storage, read_write> dst : Dst;
 
         var<workgroup> wg_struct : StructValues;
         var<workgroup> wg_matrix : mat2x2<f32>;

src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp

@@ -663,7 +663,7 @@ TEST_P(DrawIndexedIndirectTest, ValidateReusedBundleWithChangingParams) {
                 firstIndex: u32,
             }
             @group(0) @binding(0) var<uniform> input: Input;
-            @group(0) @binding(1) var<storage, write> params: Params;
+            @group(0) @binding(1) var<storage, read_write> params: Params;
             @compute @workgroup_size(1) fn main() {
                 params.indexCount = 3u;
                 params.instanceCount = 1u;

src/dawn/tests/end2end/MaxLimitTests.cpp

@@ -41,7 +41,7 @@ TEST_P(MaxLimitTests, MaxComputeWorkgroupStorageSize) {
             value1 : u32,
         }
 
-        @group(0) @binding(0) var<storage, write> dst : Dst;
+        @group(0) @binding(0) var<storage, read_write> dst : Dst;
 
         struct WGData {
           value0 : u32,
@@ -142,7 +142,7 @@ TEST_P(MaxLimitTests, MaxBufferBindingSize) {
                   }
 
                   @group(0) @binding(0) var<storage, read> buf : Buf;
-                  @group(0) @binding(1) var<storage, write> result : Result;
+                  @group(0) @binding(1) var<storage, read_write> result : Result;
 
                   @compute @workgroup_size(1,1,1)
                   fn main() {
@@ -173,7 +173,7 @@ TEST_P(MaxLimitTests, MaxBufferBindingSize) {
                   }
 
                   @group(0) @binding(0) var<uniform> buf : Buf;
-                  @group(0) @binding(1) var<storage, write> result : Result;
+                  @group(0) @binding(1) var<storage, read_write> result : Result;
 
                   @compute @workgroup_size(1,1,1)
                   fn main() {

src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp

@@ -33,7 +33,7 @@ const std::string& kMatMulFloatHeader = R"(
 
         @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(2) var<storage, read_write> resultMatrix : Matrix;
         @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 
         fn mm_readA(row : u32, col : u32) -> f32  {
@@ -200,7 +200,7 @@ const std::string& kMatMulVec4Header = R"(
 
         @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(2) var<storage, read_write> resultMatrix : Matrix;
         @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 
         fn mm_readA(row : u32, col : u32) -> vec4<f32>  {

src/tint/ast/module_clone_test.cc

@@ -52,7 +52,7 @@ var<private> g1 : f32 = 123.0;
 @group(4) @binding(0) var g6 : texture_external;
 
 var<private> g7 : vec3<f32>;
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 @group(1) @binding(1) var<storage, read> g9 : S0;
 @group(2) @binding(1) var<storage, read_write> g10 : S0;
 

src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc

@@ -351,7 +351,7 @@ struct S {
 
 var<private> a: S;
 @group(1) @binding(1) var<uniform> b: S;
-@group(1) @binding(2) var<storage, write> c: S;
+@group(1) @binding(2) var<storage, read_write> c: S;
 fn f() {
   let ptr_b = &b;
   *&a = *ptr_b;
@@ -422,7 +422,7 @@ struct S {
 };
 
 var<private> a: S;
-@group(0) @binding(0) var<storage, write> e: S;
+@group(0) @binding(0) var<storage, read_write> e: S;
 @group(1) @binding(1) var<uniform> b: S;
 fn f() {
   *&a = *&b;

src/tint/resolver/storage_class_validation_test.cc

@@ -180,6 +180,33 @@ TEST_F(ResolverStorageClassValidationTest, NotStorage_AccessMode) {
         R"(56:78 error: only variables in <storage> storage class may declare an access mode)");
 }
 
+TEST_F(ResolverStorageClassValidationTest, Storage_ReadAccessMode) {
+    // @group(0) @binding(0) var<storage, read> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kRead,
+              GroupAndBinding(0, 0));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, Storage_ReadWriteAccessMode) {
+    // @group(0) @binding(0) var<storage, read_write> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+              GroupAndBinding(0, 0));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, Storage_WriteAccessMode) {
+    // @group(0) @binding(0) var<storage, read_write> a : i32;
+    GlobalVar(Source{{56, 78}}, "a", ty.i32(), ast::StorageClass::kStorage, ast::Access::kWrite,
+              GroupAndBinding(0, 0));
+
+    ASSERT_FALSE(r()->Resolve());
+
+    EXPECT_EQ(r()->error(),
+              R"(56:78 error: access mode 'write' is not valid for the 'storage' address space)");
+}
+
 TEST_F(ResolverStorageClassValidationTest, StorageBufferNoError_Basic) {
     // struct S { x : i32 };
     // var<storage, read> g : S;

src/tint/resolver/validator.cc

@@ -554,11 +554,20 @@ bool Validator::GlobalVariable(
             // https://gpuweb.github.io/gpuweb/wgsl/#variable-declaration
             // The access mode always has a default, and except for variables in the
             // storage storage class, must not be written.
-            if (global->StorageClass() != ast::StorageClass::kStorage &&
+            if (var->declared_access != ast::Access::kUndefined) {
-                var->declared_access != ast::Access::kUndefined) {
+                if (global->StorageClass() == ast::StorageClass::kStorage) {
-                AddError("only variables in <storage> storage class may declare an access mode",
+                    // The access mode for the storage address space can only be 'read' or
-                         var->source);
+                    // 'read_write'.
-                return false;
+                    if (var->declared_access == ast::Access::kWrite) {
+                        AddError("access mode 'write' is not valid for the 'storage' address space",
+                                 decl->source);
+                        return false;
+                    }
+                } else {
+                    AddError("only variables in <storage> storage class may declare an access mode",
+                             decl->source);
+                    return false;
+                }
             }
 
             if (!AtomicVariable(global, atomic_composite_info)) {

src/tint/transform/binding_remapper_test.cc

@@ -137,9 +137,9 @@ struct S {
   a : f32,
 };
 
-@group(2) @binding(1) var<storage, read> a : S;
+@group(2) @binding(1) var<storage, read_write> a : S;
 
-@group(3) @binding(2) var<storage, write> b : S;
+@group(3) @binding(2) var<storage, read_write> b : S;
 
 @group(4) @binding(3) var<storage, read> c : S;
 
@@ -153,9 +153,9 @@ struct S {
   a : f32,
 }
 
-@group(2) @binding(1) var<storage, write> a : S;
+@group(2) @binding(1) var<storage, read_write> a : S;
 
-@group(3) @binding(2) var<storage, write> b : S;
+@group(3) @binding(2) var<storage, read_write> b : S;
 
 @group(4) @binding(3) var<storage, read> c : S;
 
@@ -168,7 +168,7 @@ fn f() {
     data.Add<BindingRemapper::Remappings>(
         BindingRemapper::BindingPoints{},
         BindingRemapper::AccessControls{
-            {{2, 1}, ast::Access::kWrite},  // Modify access control
+            {{2, 1}, ast::Access::kReadWrite},  // Modify access control
             // Keep @group(3) @binding(2) as is
             {{4, 3}, ast::Access::kRead},  // Add access control
         });
@@ -197,9 +197,9 @@ struct S {
   a : f32,
 }
 
-@group(4) @binding(5) var<storage, write> a : S;
+@group(4) @binding(5) var<storage, read_write> a : S;
 
-@group(6) @binding(7) var<storage, write> b : S;
+@group(6) @binding(7) var<storage, read_write> b : S;
 
 @compute @workgroup_size(1)
 fn f() {
@@ -213,8 +213,8 @@ fn f() {
             {{3, 2}, {6, 7}},
         },
         BindingRemapper::AccessControls{
-            {{2, 1}, ast::Access::kWrite},
+            {{2, 1}, ast::Access::kReadWrite},
-            {{3, 2}, ast::Access::kWrite},
+            {{3, 2}, ast::Access::kReadWrite},
         });
     auto got = Run<BindingRemapper>(src, data);
 

src/tint/transform/num_workgroups_from_uniform_test.cc

@@ -568,7 +568,7 @@ struct S1 {
 @group(3) @binding(0) var g5 : texture_depth_cube_array;
 @group(4) @binding(0) var g6 : texture_external;
 
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 @group(1) @binding(3) var<storage, read> g9 : S0;
 @group(3) @binding(2) var<storage, read_write> g10 : S0;
 
@@ -634,7 +634,7 @@ struct S1 {
 
 @group(4) @binding(0) var g6 : texture_external;
 
-@group(0) @binding(1) var<storage, write> g8 : S0;
+@group(0) @binding(1) var<storage, read_write> g8 : S0;
 
 @group(1) @binding(3) var<storage, read> g9 : S0;
 

src/tint/writer/glsl/generator_impl_function_test.cc

@@ -549,7 +549,7 @@ TEST_F(GlslGeneratorImplTest_Function, Emit_Attribute_EntryPoint_With_WO_Storage
                                     Member("b", ty.f32()),
                                 });
 
-    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
+    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
               ast::AttributeList{
                   create<ast::BindingAttribute>(0u),
                   create<ast::GroupAttribute>(1u),

src/tint/writer/hlsl/generator_impl_function_test.cc

@@ -503,7 +503,7 @@ TEST_F(HlslGeneratorImplTest_Function, Emit_Attribute_EntryPoint_With_WO_Storage
                                     Member("b", ty.f32()),
                                 });
 
-    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
+    GlobalVar("coord", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
               ast::AttributeList{
                   create<ast::BindingAttribute>(0u),
                   create<ast::GroupAttribute>(1u),

src/tint/writer/wgsl/generator_impl_variable_test.cc

@@ -56,21 +56,6 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Read) {
     EXPECT_EQ(out.str(), R"(@binding(0) @group(0) var<storage, read> a : S;)");
 }
 
-TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Write) {
-    auto* s = Structure("S", {Member("a", ty.i32())});
-    auto* v = GlobalVar("a", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kWrite,
-                        ast::AttributeList{
-                            create<ast::BindingAttribute>(0u),
-                            create<ast::GroupAttribute>(0u),
-                        });
-
-    GeneratorImpl& gen = Build();
-
-    std::stringstream out;
-    ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
-    EXPECT_EQ(out.str(), R"(@binding(0) @group(0) var<storage, write> a : S;)");
-}
-
 TEST_F(WgslGeneratorImplTest, EmitVariable_Access_ReadWrite) {
     auto* s = Structure("S", {Member("a", ty.i32())});
     auto* v = GlobalVar("a", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,

test/tint/benchmark/metaball-isosurface.wgsl

@@ -14,25 +14,25 @@ struct IsosurfaceVolume {
   values : array<f32>,
 }
 
-@group(0) @binding(1) var<storage, write> volume : IsosurfaceVolume;
+@group(0) @binding(1) var<storage, read_write> volume : IsosurfaceVolume;
 
 struct PositionBuffer {
   values : array<f32>,
 }
 
-@group(0) @binding(2) var<storage, write> positionsOut : PositionBuffer;
+@group(0) @binding(2) var<storage, read_write> positionsOut : PositionBuffer;
 
 struct NormalBuffer {
   values : array<f32>,
 }
 
-@group(0) @binding(3) var<storage, write> normalsOut : NormalBuffer;
+@group(0) @binding(3) var<storage, read_write> normalsOut : NormalBuffer;
 
 struct IndexBuffer {
   tris : array<u32>,
 }
 
-@group(0) @binding(4) var<storage, write> indicesOut : IndexBuffer;
+@group(0) @binding(4) var<storage, read_write> indicesOut : IndexBuffer;
 
 struct DrawIndirectArgs {
   vc : u32,

test/tint/buffer/storage/static_index/write.wgsl

@@ -15,7 +15,7 @@ struct S {
     j : array<Inner, 4>,
 };
 
-@binding(0) @group(0) var<storage, write> s : S;
+@binding(0) @group(0) var<storage, read_write> s : S;
 
 @compute @workgroup_size(1)
 fn main() {

test/tint/buffer/storage/static_index/write.wgsl.expected.spvasm

@@ -39,7 +39,6 @@
                OpMemberDecorate %Inner 0 Offset 0
                OpMemberDecorate %S 9 Offset 108
                OpDecorate %_arr_Inner_uint_4 ArrayStride 4
-               OpDecorate %s NonReadable
                OpDecorate %s Binding 0
                OpDecorate %s DescriptorSet 0
         %int = OpTypeInt 32 1

test/tint/buffer/storage/static_index/write.wgsl.expected.wgsl

@@ -15,7 +15,7 @@ struct S {
   j : array<Inner, 4>,
 }
 
-@binding(0) @group(0) var<storage, write> s : S;
+@binding(0) @group(0) var<storage, read_write> s : S;
 
 @compute @workgroup_size(1)
 fn main() {

test/tint/bug/tint/744.wgsl

@@ -9,7 +9,7 @@
 
 @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(2) var<storage, read_write> resultMatrix : Matrix;
 @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 
 @compute @workgroup_size(2,2,1)

test/tint/bug/tint/744.wgsl.expected.spvasm

@@ -33,7 +33,6 @@
                OpDecorate %secondMatrix NonWritable
                OpDecorate %secondMatrix DescriptorSet 0
                OpDecorate %secondMatrix Binding 1
-               OpDecorate %resultMatrix NonReadable
                OpDecorate %resultMatrix DescriptorSet 0
                OpDecorate %resultMatrix Binding 2
                OpDecorate %Uniforms Block

test/tint/bug/tint/744.wgsl.expected.wgsl

@@ -12,7 +12,7 @@ struct Matrix {
 
 @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
 
-@group(0) @binding(2) var<storage, write> resultMatrix : Matrix;
+@group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;
 
 @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 

test/tint/bug/tint/914.wgsl

@@ -9,7 +9,7 @@
 
 @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(2) var<storage, read_write> resultMatrix : Matrix;
 @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 
 fn mm_readA(row : u32, col : u32) -> f32  {

test/tint/bug/tint/914.wgsl.expected.spvasm

@@ -65,7 +65,6 @@
                OpDecorate %secondMatrix NonWritable
                OpDecorate %secondMatrix DescriptorSet 0
                OpDecorate %secondMatrix Binding 1
-               OpDecorate %resultMatrix NonReadable
                OpDecorate %resultMatrix DescriptorSet 0
                OpDecorate %resultMatrix Binding 2
                OpDecorate %Uniforms Block

test/tint/bug/tint/914.wgsl.expected.wgsl

@@ -12,7 +12,7 @@ struct Matrix {
 
 @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
 
-@group(0) @binding(2) var<storage, write> resultMatrix : Matrix;
+@group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;
 
 @group(0) @binding(3) var<uniform> uniforms : Uniforms;
 

test/tint/bug/tint/993.wgsl

@@ -7,7 +7,7 @@
  struct Result {
   value: u32,
 };
-@group(1) @binding(1) var<storage, write> result: Result;
+@group(1) @binding(1) var<storage, read_write> result: Result;
 
  struct TestData {
   data: array<atomic<i32>,3>,

test/tint/bug/tint/993.wgsl.expected.spvasm

@@ -25,7 +25,6 @@
                OpDecorate %constants Binding 0
                OpDecorate %Result Block
                OpMemberDecorate %Result 0 Offset 0
-               OpDecorate %result NonReadable
                OpDecorate %result DescriptorSet 1
                OpDecorate %result Binding 1
                OpDecorate %TestData Block

test/tint/bug/tint/993.wgsl.expected.wgsl

@@ -8,7 +8,7 @@ struct Result {
   value : u32,
 }
 
-@group(1) @binding(1) var<storage, write> result : Result;
+@group(1) @binding(1) var<storage, read_write> result : Result;
 
 struct TestData {
   data : array<atomic<i32>, 3>,

test/tint/bug/tint/998.wgsl

@@ -6,7 +6,7 @@
  struct Result {
   value: u32,
 };
-@group(1) @binding(1) var<storage, write> result: Result;
+@group(1) @binding(1) var<storage, read_write> result: Result;
 
 struct S {
   data: array<u32, 3>,

test/tint/bug/tint/998.wgsl.expected.spvasm

@@ -24,7 +24,6 @@
                OpDecorate %constants Binding 0
                OpDecorate %Result Block
                OpMemberDecorate %Result 0 Offset 0
-               OpDecorate %result NonReadable
                OpDecorate %result DescriptorSet 1
                OpDecorate %result Binding 1
                OpMemberDecorate %S 0 Offset 0

test/tint/bug/tint/998.wgsl.expected.wgsl

@@ -8,7 +8,7 @@ struct Result {
   value : u32,
 }
 
-@group(1) @binding(1) var<storage, write> result : Result;
+@group(1) @binding(1) var<storage, read_write> result : Result;
 
 struct S {
   data : array<u32, 3>,

test/tint/shader_io/shared_struct_storage_buffer.wgsl

@@ -5,7 +5,7 @@ struct S {
 };
 
 @group(0) @binding(0)
-var<storage, write> output : S;
+var<storage, read_write> output : S;
 
 @fragment
 fn frag_main(input : S) {

test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.spvasm

@@ -26,7 +26,6 @@
                OpMemberDecorate %S 0 Offset 0
                OpMemberDecorate %S 1 Offset 4
                OpMemberDecorate %S 2 Offset 128
-               OpDecorate %output NonReadable
                OpDecorate %output DescriptorSet 0
                OpDecorate %output Binding 0
       %float = OpTypeFloat 32

test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.wgsl

@@ -7,7 +7,7 @@ struct S {
   v : vec4<f32>,
 }
 
-@group(0) @binding(0) var<storage, write> output : S;
+@group(0) @binding(0) var<storage, read_write> output : S;
 
 @fragment
 fn frag_main(input : S) {