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>
diff --git a/src/dawn/native/ComputePassEncoder.cpp b/src/dawn/native/ComputePassEncoder.cpp
index e70aea3..374ff6f 100644
--- a/src/dawn/native/ComputePassEncoder.cpp
+++ b/src/dawn/native/ComputePassEncoder.cpp
@@ -62,7 +62,7 @@
@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() {
diff --git a/src/dawn/native/IndirectDrawValidationEncoder.cpp b/src/dawn/native/IndirectDrawValidationEncoder.cpp
index 10a8164..abd09cd 100644
--- a/src/dawn/native/IndirectDrawValidationEncoder.cpp
+++ b/src/dawn/native/IndirectDrawValidationEncoder.cpp
@@ -81,7 +81,7 @@
@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;
diff --git a/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp b/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp
index f69b5df..c63e762 100644
--- a/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp
+++ b/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp
@@ -78,7 +78,7 @@
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 @@
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>;
diff --git a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
index 78ee799..9e99598 100644
--- a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
+++ b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp
@@ -663,7 +663,7 @@
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;
diff --git a/src/dawn/tests/end2end/MaxLimitTests.cpp b/src/dawn/tests/end2end/MaxLimitTests.cpp
index fa17ab3..052713f 100644
--- a/src/dawn/tests/end2end/MaxLimitTests.cpp
+++ b/src/dawn/tests/end2end/MaxLimitTests.cpp
@@ -41,7 +41,7 @@
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 @@
}
@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 @@
}
@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() {
diff --git a/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp b/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp
index d3faea0..5a68e2e 100644
--- a/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp
+++ b/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp
@@ -33,7 +33,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 {
@@ -200,7 +200,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) -> vec4<f32> {
diff --git a/src/tint/ast/module_clone_test.cc b/src/tint/ast/module_clone_test.cc
index 544e6bf..a79ef0e 100644
--- a/src/tint/ast/module_clone_test.cc
+++ b/src/tint/ast/module_clone_test.cc
@@ -52,7 +52,7 @@
@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;
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
index 737cb07..c2d40d9 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
@@ -351,7 +351,7 @@
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 @@
};
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;
diff --git a/src/tint/resolver/storage_class_validation_test.cc b/src/tint/resolver/storage_class_validation_test.cc
index 0e2be29..3b98fdb 100644
--- a/src/tint/resolver/storage_class_validation_test.cc
+++ b/src/tint/resolver/storage_class_validation_test.cc
@@ -180,6 +180,33 @@
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;
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index a2e7799..ee98d03 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -554,11 +554,20 @@
// 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 &&
- var->declared_access != ast::Access::kUndefined) {
- AddError("only variables in <storage> storage class may declare an access mode",
- var->source);
- return false;
+ if (var->declared_access != ast::Access::kUndefined) {
+ if (global->StorageClass() == ast::StorageClass::kStorage) {
+ // The access mode for the storage address space can only be 'read' or
+ // 'read_write'.
+ 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)) {
diff --git a/src/tint/transform/binding_remapper_test.cc b/src/tint/transform/binding_remapper_test.cc
index 3274886..564a3a5 100644
--- a/src/tint/transform/binding_remapper_test.cc
+++ b/src/tint/transform/binding_remapper_test.cc
@@ -137,9 +137,9 @@
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 @@
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 @@
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 @@
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 @@
{{3, 2}, {6, 7}},
},
BindingRemapper::AccessControls{
- {{2, 1}, ast::Access::kWrite},
- {{3, 2}, ast::Access::kWrite},
+ {{2, 1}, ast::Access::kReadWrite},
+ {{3, 2}, ast::Access::kReadWrite},
});
auto got = Run<BindingRemapper>(src, data);
diff --git a/src/tint/transform/num_workgroups_from_uniform_test.cc b/src/tint/transform/num_workgroups_from_uniform_test.cc
index 8562c01..093081c 100644
--- a/src/tint/transform/num_workgroups_from_uniform_test.cc
+++ b/src/tint/transform/num_workgroups_from_uniform_test.cc
@@ -568,7 +568,7 @@
@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 @@
@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;
diff --git a/src/tint/writer/glsl/generator_impl_function_test.cc b/src/tint/writer/glsl/generator_impl_function_test.cc
index a70e238..450041b 100644
--- a/src/tint/writer/glsl/generator_impl_function_test.cc
+++ b/src/tint/writer/glsl/generator_impl_function_test.cc
@@ -549,7 +549,7 @@
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),
diff --git a/src/tint/writer/hlsl/generator_impl_function_test.cc b/src/tint/writer/hlsl/generator_impl_function_test.cc
index 85647a5..89d74c5 100644
--- a/src/tint/writer/hlsl/generator_impl_function_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_function_test.cc
@@ -503,7 +503,7 @@
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),
diff --git a/src/tint/writer/wgsl/generator_impl_variable_test.cc b/src/tint/writer/wgsl/generator_impl_variable_test.cc
index a02aed8..a058bc7 100644
--- a/src/tint/writer/wgsl/generator_impl_variable_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_variable_test.cc
@@ -56,21 +56,6 @@
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,
diff --git a/test/tint/benchmark/metaball-isosurface.wgsl b/test/tint/benchmark/metaball-isosurface.wgsl
index 33eecb0..beecb2f 100644
--- a/test/tint/benchmark/metaball-isosurface.wgsl
+++ b/test/tint/benchmark/metaball-isosurface.wgsl
@@ -14,25 +14,25 @@
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,
diff --git a/test/tint/buffer/storage/static_index/write.wgsl b/test/tint/buffer/storage/static_index/write.wgsl
index e1f7eab..5290db8 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl
+++ b/test/tint/buffer/storage/static_index/write.wgsl
@@ -15,7 +15,7 @@
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() {
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.spvasm b/test/tint/buffer/storage/static_index/write.wgsl.expected.spvasm
index c099dbd..5f06f4e 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.wgsl b/test/tint/buffer/storage/static_index/write.wgsl.expected.wgsl
index fd7a137..af8c2de 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.wgsl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.wgsl
@@ -15,7 +15,7 @@
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() {
diff --git a/test/tint/bug/tint/744.wgsl b/test/tint/bug/tint/744.wgsl
index 70e57cc..a978f81 100644
--- a/test/tint/bug/tint/744.wgsl
+++ b/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)
diff --git a/test/tint/bug/tint/744.wgsl.expected.spvasm b/test/tint/bug/tint/744.wgsl.expected.spvasm
index 2bc71b4..f83372b 100644
--- a/test/tint/bug/tint/744.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/bug/tint/744.wgsl.expected.wgsl b/test/tint/bug/tint/744.wgsl.expected.wgsl
index d08d2d9..ec33a80 100644
--- a/test/tint/bug/tint/744.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/744.wgsl.expected.wgsl
@@ -12,7 +12,7 @@
@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;
diff --git a/test/tint/bug/tint/914.wgsl b/test/tint/bug/tint/914.wgsl
index 6c0d6c3..1fde262 100644
--- a/test/tint/bug/tint/914.wgsl
+++ b/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 {
diff --git a/test/tint/bug/tint/914.wgsl.expected.spvasm b/test/tint/bug/tint/914.wgsl.expected.spvasm
index 8647c10..6506377 100644
--- a/test/tint/bug/tint/914.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/bug/tint/914.wgsl.expected.wgsl b/test/tint/bug/tint/914.wgsl.expected.wgsl
index 7b1cedb..10301a0 100644
--- a/test/tint/bug/tint/914.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/914.wgsl.expected.wgsl
@@ -12,7 +12,7 @@
@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;
diff --git a/test/tint/bug/tint/993.wgsl b/test/tint/bug/tint/993.wgsl
index e45d789..345c3cc 100644
--- a/test/tint/bug/tint/993.wgsl
+++ b/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>,
diff --git a/test/tint/bug/tint/993.wgsl.expected.spvasm b/test/tint/bug/tint/993.wgsl.expected.spvasm
index 07e03b2..9abd51a 100644
--- a/test/tint/bug/tint/993.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/bug/tint/993.wgsl.expected.wgsl b/test/tint/bug/tint/993.wgsl.expected.wgsl
index df926b8..ef33e7c 100644
--- a/test/tint/bug/tint/993.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/993.wgsl.expected.wgsl
@@ -8,7 +8,7 @@
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>,
diff --git a/test/tint/bug/tint/998.wgsl b/test/tint/bug/tint/998.wgsl
index e04eec0..139e461 100644
--- a/test/tint/bug/tint/998.wgsl
+++ b/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>,
diff --git a/test/tint/bug/tint/998.wgsl.expected.spvasm b/test/tint/bug/tint/998.wgsl.expected.spvasm
index 8d3cc74..82231fb 100644
--- a/test/tint/bug/tint/998.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/bug/tint/998.wgsl.expected.wgsl b/test/tint/bug/tint/998.wgsl.expected.wgsl
index ebdae6e..e05a05c 100644
--- a/test/tint/bug/tint/998.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/998.wgsl.expected.wgsl
@@ -8,7 +8,7 @@
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>,
diff --git a/test/tint/shader_io/shared_struct_storage_buffer.wgsl b/test/tint/shader_io/shared_struct_storage_buffer.wgsl
index d5f239a..2d84447 100644
--- a/test/tint/shader_io/shared_struct_storage_buffer.wgsl
+++ b/test/tint/shader_io/shared_struct_storage_buffer.wgsl
@@ -5,7 +5,7 @@
};
@group(0) @binding(0)
-var<storage, write> output : S;
+var<storage, read_write> output : S;
@fragment
fn frag_main(input : S) {
diff --git a/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.spvasm b/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.spvasm
index 5f96817..e2748bd 100644
--- a/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.spvasm
+++ b/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
diff --git a/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.wgsl b/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.wgsl
index 38fe6ef..d12912e 100644
--- a/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.wgsl
+++ b/test/tint/shader_io/shared_struct_storage_buffer.wgsl.expected.wgsl
@@ -7,7 +7,7 @@
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) {
