From 5286ea9d168aa5a38421d2ee6febe6abcdf0ded3 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Mon, 4 Jul 2022 15:17:00 +0000 Subject: [PATCH] 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 Commit-Queue: Dan Sinclair Reviewed-by: Dan Sinclair --- src/dawn/native/ComputePassEncoder.cpp | 2 +- .../native/IndirectDrawValidationEncoder.cpp | 2 +- .../end2end/ComputeSharedMemoryTests.cpp | 4 +-- .../end2end/DrawIndexedIndirectTests.cpp | 2 +- src/dawn/tests/end2end/MaxLimitTests.cpp | 6 ++--- .../tests/perf_tests/ShaderRobustnessPerf.cpp | 4 +-- src/tint/ast/module_clone_test.cc | 2 +- .../mutations/replace_identifier_test.cc | 4 +-- .../resolver/storage_class_validation_test.cc | 27 +++++++++++++++++++ src/tint/resolver/validator.cc | 19 +++++++++---- src/tint/transform/binding_remapper_test.cc | 18 ++++++------- .../num_workgroups_from_uniform_test.cc | 4 +-- .../glsl/generator_impl_function_test.cc | 2 +- .../hlsl/generator_impl_function_test.cc | 2 +- .../wgsl/generator_impl_variable_test.cc | 15 ----------- test/tint/benchmark/metaball-isosurface.wgsl | 8 +++--- .../buffer/storage/static_index/write.wgsl | 2 +- .../static_index/write.wgsl.expected.spvasm | 1 - .../static_index/write.wgsl.expected.wgsl | 2 +- test/tint/bug/tint/744.wgsl | 2 +- test/tint/bug/tint/744.wgsl.expected.spvasm | 1 - test/tint/bug/tint/744.wgsl.expected.wgsl | 2 +- test/tint/bug/tint/914.wgsl | 2 +- test/tint/bug/tint/914.wgsl.expected.spvasm | 1 - test/tint/bug/tint/914.wgsl.expected.wgsl | 2 +- test/tint/bug/tint/993.wgsl | 2 +- test/tint/bug/tint/993.wgsl.expected.spvasm | 1 - test/tint/bug/tint/993.wgsl.expected.wgsl | 2 +- test/tint/bug/tint/998.wgsl | 2 +- test/tint/bug/tint/998.wgsl.expected.spvasm | 1 - test/tint/bug/tint/998.wgsl.expected.wgsl | 2 +- .../shared_struct_storage_buffer.wgsl | 2 +- ...struct_storage_buffer.wgsl.expected.spvasm | 1 - ...d_struct_storage_buffer.wgsl.expected.wgsl | 2 +- 34 files changed, 83 insertions(+), 68 deletions(-) diff --git a/src/dawn/native/ComputePassEncoder.cpp b/src/dawn/native/ComputePassEncoder.cpp index e70aea3193..374ff6f483 100644 --- a/src/dawn/native/ComputePassEncoder.cpp +++ b/src/dawn/native/ComputePassEncoder.cpp @@ -62,7 +62,7 @@ ResultOrError GetOrCreateIndirectDispatchValidationPipelin @group(0) @binding(0) var uniformParams: UniformParams; @group(0) @binding(1) var clientParams: IndirectParams; - @group(0) @binding(2) var validatedParams: ValidatedParams; + @group(0) @binding(2) var 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 10a8164026..abd09cd112 100644 --- a/src/dawn/native/IndirectDrawValidationEncoder.cpp +++ b/src/dawn/native/IndirectDrawValidationEncoder.cpp @@ -81,7 +81,7 @@ static const char sRenderValidationShaderSource[] = R"( @group(0) @binding(0) var batch: BatchInfo; @group(0) @binding(1) var inputParams: IndirectParams; - @group(0) @binding(2) var outputParams: IndirectParams; + @group(0) @binding(2) var 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 f69b5df845..c63e762285 100644 --- a/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp +++ b/src/dawn/tests/end2end/ComputeSharedMemoryTests.cpp @@ -78,7 +78,7 @@ TEST_P(ComputeSharedMemoryTests, Basic) { x : u32 } - @group(0) @binding(0) var dst : Dst; + @group(0) @binding(0) var dst : Dst; var tmp : u32; @compute @workgroup_size(4,4,1) @@ -117,7 +117,7 @@ TEST_P(ComputeSharedMemoryTests, AssortedTypes) { d_vector : vec4, } - @group(0) @binding(0) var dst : Dst; + @group(0) @binding(0) var dst : Dst; var wg_struct : StructValues; var wg_matrix : mat2x2; diff --git a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp index 78ee799566..9e995988e4 100644 --- a/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp +++ b/src/dawn/tests/end2end/DrawIndexedIndirectTests.cpp @@ -663,7 +663,7 @@ TEST_P(DrawIndexedIndirectTest, ValidateReusedBundleWithChangingParams) { firstIndex: u32, } @group(0) @binding(0) var input: Input; - @group(0) @binding(1) var params: Params; + @group(0) @binding(1) var 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 fa17ab3eb9..052713fb96 100644 --- a/src/dawn/tests/end2end/MaxLimitTests.cpp +++ b/src/dawn/tests/end2end/MaxLimitTests.cpp @@ -41,7 +41,7 @@ TEST_P(MaxLimitTests, MaxComputeWorkgroupStorageSize) { value1 : u32, } - @group(0) @binding(0) var dst : Dst; + @group(0) @binding(0) var dst : Dst; struct WGData { value0 : u32, @@ -142,7 +142,7 @@ TEST_P(MaxLimitTests, MaxBufferBindingSize) { } @group(0) @binding(0) var buf : Buf; - @group(0) @binding(1) var result : Result; + @group(0) @binding(1) var result : Result; @compute @workgroup_size(1,1,1) fn main() { @@ -173,7 +173,7 @@ TEST_P(MaxLimitTests, MaxBufferBindingSize) { } @group(0) @binding(0) var buf : Buf; - @group(0) @binding(1) var result : Result; + @group(0) @binding(1) var 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 d3faea0aea..5a68e2e3b9 100644 --- a/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp +++ b/src/dawn/tests/perf_tests/ShaderRobustnessPerf.cpp @@ -33,7 +33,7 @@ const std::string& kMatMulFloatHeader = R"( @group(0) @binding(0) var firstMatrix : Matrix; @group(0) @binding(1) var secondMatrix : Matrix; - @group(0) @binding(2) var resultMatrix : Matrix; + @group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var uniforms : Uniforms; fn mm_readA(row : u32, col : u32) -> f32 { @@ -200,7 +200,7 @@ const std::string& kMatMulVec4Header = R"( @group(0) @binding(0) var firstMatrix : Matrix; @group(0) @binding(1) var secondMatrix : Matrix; - @group(0) @binding(2) var resultMatrix : Matrix; + @group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var uniforms : Uniforms; fn mm_readA(row : u32, col : u32) -> vec4 { diff --git a/src/tint/ast/module_clone_test.cc b/src/tint/ast/module_clone_test.cc index 544e6bf323..a79ef0e566 100644 --- a/src/tint/ast/module_clone_test.cc +++ b/src/tint/ast/module_clone_test.cc @@ -52,7 +52,7 @@ var g1 : f32 = 123.0; @group(4) @binding(0) var g6 : texture_external; var g7 : vec3; -@group(0) @binding(1) var g8 : S0; +@group(0) @binding(1) var g8 : S0; @group(1) @binding(1) var g9 : S0; @group(2) @binding(1) var 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 737cb07c2e..c2d40d9a78 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 @@ struct S { var a: S; @group(1) @binding(1) var b: S; -@group(1) @binding(2) var c: S; +@group(1) @binding(2) var c: S; fn f() { let ptr_b = &b; *&a = *ptr_b; @@ -422,7 +422,7 @@ struct S { }; var a: S; -@group(0) @binding(0) var e: S; +@group(0) @binding(0) var e: S; @group(1) @binding(1) var 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 0e2be29d18..3b98fdb5d5 100644 --- a/src/tint/resolver/storage_class_validation_test.cc +++ 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 class may declare an access mode)"); } +TEST_F(ResolverStorageClassValidationTest, Storage_ReadAccessMode) { + // @group(0) @binding(0) var 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 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 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 g : S; diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index a2e7799887..ee98d03ac9 100644 --- a/src/tint/resolver/validator.cc +++ b/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 && - var->declared_access != ast::Access::kUndefined) { - AddError("only variables in 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 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 3274886608..564a3a5b24 100644 --- a/src/tint/transform/binding_remapper_test.cc +++ b/src/tint/transform/binding_remapper_test.cc @@ -137,9 +137,9 @@ struct S { a : f32, }; -@group(2) @binding(1) var a : S; +@group(2) @binding(1) var a : S; -@group(3) @binding(2) var b : S; +@group(3) @binding(2) var b : S; @group(4) @binding(3) var c : S; @@ -153,9 +153,9 @@ struct S { a : f32, } -@group(2) @binding(1) var a : S; +@group(2) @binding(1) var a : S; -@group(3) @binding(2) var b : S; +@group(3) @binding(2) var b : S; @group(4) @binding(3) var c : S; @@ -168,7 +168,7 @@ fn f() { data.Add( 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 a : S; +@group(4) @binding(5) var a : S; -@group(6) @binding(7) var b : S; +@group(6) @binding(7) var 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}, - {{3, 2}, ast::Access::kWrite}, + {{2, 1}, ast::Access::kReadWrite}, + {{3, 2}, ast::Access::kReadWrite}, }); auto got = Run(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 8562c01d2a..093081c6e7 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 @@ 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 g8 : S0; +@group(0) @binding(1) var g8 : S0; @group(1) @binding(3) var g9 : S0; @group(3) @binding(2) var g10 : S0; @@ -634,7 +634,7 @@ struct S1 { @group(4) @binding(0) var g6 : texture_external; -@group(0) @binding(1) var g8 : S0; +@group(0) @binding(1) var g8 : S0; @group(1) @binding(3) var 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 a70e238e8f..450041b8a1 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 @@ 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(0u), create(1u), diff --git a/src/tint/writer/hlsl/generator_impl_function_test.cc b/src/tint/writer/hlsl/generator_impl_function_test.cc index 85647a5b45..89d74c54dd 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 @@ 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(0u), create(1u), diff --git a/src/tint/writer/wgsl/generator_impl_variable_test.cc b/src/tint/writer/wgsl/generator_impl_variable_test.cc index a02aed8789..a058bc7cef 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 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Read) { EXPECT_EQ(out.str(), R"(@binding(0) @group(0) var 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(0u), - create(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 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 33eecb041f..beecb2f42c 100644 --- a/test/tint/benchmark/metaball-isosurface.wgsl +++ b/test/tint/benchmark/metaball-isosurface.wgsl @@ -14,25 +14,25 @@ struct IsosurfaceVolume { values : array, } -@group(0) @binding(1) var volume : IsosurfaceVolume; +@group(0) @binding(1) var volume : IsosurfaceVolume; struct PositionBuffer { values : array, } -@group(0) @binding(2) var positionsOut : PositionBuffer; +@group(0) @binding(2) var positionsOut : PositionBuffer; struct NormalBuffer { values : array, } -@group(0) @binding(3) var normalsOut : NormalBuffer; +@group(0) @binding(3) var normalsOut : NormalBuffer; struct IndexBuffer { tris : array, } -@group(0) @binding(4) var indicesOut : IndexBuffer; +@group(0) @binding(4) var 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 e1f7eab7e3..5290db8051 100644 --- a/test/tint/buffer/storage/static_index/write.wgsl +++ b/test/tint/buffer/storage/static_index/write.wgsl @@ -15,7 +15,7 @@ struct S { j : array, }; -@binding(0) @group(0) var s : S; +@binding(0) @group(0) var 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 c099dbda02..5f06f4ef83 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 fd7a1376a4..af8c2de8d4 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 @@ struct S { j : array, } -@binding(0) @group(0) var s : S; +@binding(0) @group(0) var 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 70e57cc605..a978f81573 100644 --- a/test/tint/bug/tint/744.wgsl +++ b/test/tint/bug/tint/744.wgsl @@ -9,7 +9,7 @@ @group(0) @binding(0) var firstMatrix : Matrix; @group(0) @binding(1) var secondMatrix : Matrix; -@group(0) @binding(2) var resultMatrix : Matrix; +@group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var 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 2bc71b48d9..f83372b822 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 d08d2d9b39..ec33a80731 100644 --- a/test/tint/bug/tint/744.wgsl.expected.wgsl +++ b/test/tint/bug/tint/744.wgsl.expected.wgsl @@ -12,7 +12,7 @@ struct Matrix { @group(0) @binding(1) var secondMatrix : Matrix; -@group(0) @binding(2) var resultMatrix : Matrix; +@group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var uniforms : Uniforms; diff --git a/test/tint/bug/tint/914.wgsl b/test/tint/bug/tint/914.wgsl index 6c0d6c398d..1fde2626a8 100644 --- a/test/tint/bug/tint/914.wgsl +++ b/test/tint/bug/tint/914.wgsl @@ -9,7 +9,7 @@ @group(0) @binding(0) var firstMatrix : Matrix; @group(0) @binding(1) var secondMatrix : Matrix; -@group(0) @binding(2) var resultMatrix : Matrix; +@group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var 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 8647c10d01..6506377514 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 7b1cedb765..10301a011e 100644 --- a/test/tint/bug/tint/914.wgsl.expected.wgsl +++ b/test/tint/bug/tint/914.wgsl.expected.wgsl @@ -12,7 +12,7 @@ struct Matrix { @group(0) @binding(1) var secondMatrix : Matrix; -@group(0) @binding(2) var resultMatrix : Matrix; +@group(0) @binding(2) var resultMatrix : Matrix; @group(0) @binding(3) var uniforms : Uniforms; diff --git a/test/tint/bug/tint/993.wgsl b/test/tint/bug/tint/993.wgsl index e45d78982b..345c3cce82 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 result: Result; +@group(1) @binding(1) var result: Result; struct TestData { data: array,3>, diff --git a/test/tint/bug/tint/993.wgsl.expected.spvasm b/test/tint/bug/tint/993.wgsl.expected.spvasm index 07e03b21fc..9abd51a6db 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 df926b80cf..ef33e7ccc6 100644 --- a/test/tint/bug/tint/993.wgsl.expected.wgsl +++ b/test/tint/bug/tint/993.wgsl.expected.wgsl @@ -8,7 +8,7 @@ struct Result { value : u32, } -@group(1) @binding(1) var result : Result; +@group(1) @binding(1) var result : Result; struct TestData { data : array, 3>, diff --git a/test/tint/bug/tint/998.wgsl b/test/tint/bug/tint/998.wgsl index e04eec0c60..139e461882 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 result: Result; +@group(1) @binding(1) var result: Result; struct S { data: array, diff --git a/test/tint/bug/tint/998.wgsl.expected.spvasm b/test/tint/bug/tint/998.wgsl.expected.spvasm index 8d3cc7422a..82231fbc02 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 ebdae6e633..e05a05cdf4 100644 --- a/test/tint/bug/tint/998.wgsl.expected.wgsl +++ b/test/tint/bug/tint/998.wgsl.expected.wgsl @@ -8,7 +8,7 @@ struct Result { value : u32, } -@group(1) @binding(1) var result : Result; +@group(1) @binding(1) var result : Result; struct S { data : array, diff --git a/test/tint/shader_io/shared_struct_storage_buffer.wgsl b/test/tint/shader_io/shared_struct_storage_buffer.wgsl index d5f239a620..2d8444767e 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 @@ struct S { }; @group(0) @binding(0) -var output : S; +var 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 5f96817e52..e2748bdaa0 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 38fe6ef54d..d12912e629 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 @@ struct S { v : vec4, } -@group(0) @binding(0) var output : S; +@group(0) @binding(0) var output : S; @fragment fn frag_main(input : S) {