wgsl: Clean up duplicated [[block]] structs

crbug.com/tint/386 has been fixed.
We don't need these work arounds any more.

Original CL by bclayton@:
https://dawn-review.googlesource.com/c/dawn/+/54641

Bug: tint:386
Fixed: dawn:943
Fixed: dawn:975
Change-Id: I1b94d1e33b023b7d9b3153094f829885a75b5a5e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/57743
Commit-Queue: Austin Eng <enga@chromium.org>
Auto-Submit: Austin Eng <enga@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/tests/end2end/BindGroupTests.cpp

@@ -338,18 +338,12 @@ TEST_P(BindGroupTests, MultipleBindLayouts) {
 
     wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
         // TODO(crbug.com/tint/369): Use a mat2x2 when Tint translates it correctly.
-        // TODO(crbug.com/tint/386): Use the same struct.
+        [[block]] struct VertexUniformBuffer {
-        [[block]] struct VertexUniformBuffer1 {
             transform : vec4<f32>;
         };
 
-        [[block]] struct VertexUniformBuffer2 {
+        [[group(0), binding(0)]] var <uniform> vertexUbo1 : VertexUniformBuffer;
-            transform : vec4<f32>;
+        [[group(1), binding(0)]] var <uniform> vertexUbo2 : VertexUniformBuffer;
-        };
-
-        // TODO(crbug.com/tint/386): Use the same struct definition.
-        [[group(0), binding(0)]] var <uniform> vertexUbo1 : VertexUniformBuffer1;
-        [[group(1), binding(0)]] var <uniform> vertexUbo2 : VertexUniformBuffer2;
 
         [[stage(vertex)]]
         fn main([[builtin(vertex_index)]] VertexIndex : u32) -> [[builtin(position)]] vec4<f32> {
@@ -365,18 +359,12 @@ TEST_P(BindGroupTests, MultipleBindLayouts) {
         })");
 
     wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
-        // TODO(crbug.com/tint/386): Use the same struct
+        [[block]] struct FragmentUniformBuffer {
-        [[block]] struct FragmentUniformBuffer1 {
             color : vec4<f32>;
         };
 
-        [[block]] struct FragmentUniformBuffer2 {
+        [[group(0), binding(1)]] var <uniform> fragmentUbo1 : FragmentUniformBuffer;
-            color : vec4<f32>;
+        [[group(1), binding(1)]] var <uniform> fragmentUbo2 : FragmentUniformBuffer;
-        };
-
-        // TODO(crbug.com/tint/386): Use the same struct definition.
-        [[group(0), binding(1)]] var <uniform> fragmentUbo1 : FragmentUniformBuffer1;
-        [[group(1), binding(1)]] var <uniform> fragmentUbo2 : FragmentUniformBuffer2;
 
         [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
             return fragmentUbo1.color + fragmentUbo2.color;
@@ -770,6 +758,9 @@ TEST_P(BindGroupTests, DrawThenChangePipelineAndBindGroup) {
 // Regression test for crbug.com/dawn/408 where dynamic offsets were applied in the wrong order.
 // Dynamic offsets should be applied in increasing order of binding number.
 TEST_P(BindGroupTests, DynamicOffsetOrder) {
+    // Does not work with SPIRV-Cross. crbug.com/dawn/975
+    DAWN_SUPPRESS_TEST_IF(IsD3D12() && !HasToggleEnabled("use_tint_generator"));
+
     // We will put the following values and the respective offsets into a buffer.
     // The test will ensure that the correct dynamic offset is applied to each buffer by reading the
     // value from an offset binding.
@@ -819,16 +810,7 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
 
     wgpu::ComputePipelineDescriptor pipelineDescriptor;
     pipelineDescriptor.compute.module = utils::CreateShaderModule(device, R"(
-        // TODO(crbug.com/tint/386): Use the same struct
+        [[block]] struct Buffer {
-        [[block]] struct Buffer0 {
-            value : u32;
-        };
-
-        [[block]] struct Buffer2 {
-            value : u32;
-        };
-
-        [[block]] struct Buffer3 {
             value : u32;
         };
 
@@ -836,9 +818,9 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
             value : vec3<u32>;
         };
 
-        [[group(0), binding(2)]] var<uniform> buffer2 : Buffer2;
+        [[group(0), binding(2)]] var<uniform> buffer2 : Buffer;
-        [[group(0), binding(3)]] var<storage, read> buffer3 : Buffer3;
+        [[group(0), binding(3)]] var<storage, read> buffer3 : Buffer;
-        [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0;
+        [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer;
         [[group(0), binding(4)]] var<storage, read_write> outputBuffer : OutputBuffer;
 
         [[stage(compute), workgroup_size(1)]] fn main() {
@@ -948,23 +930,13 @@ TEST_P(BindGroupTests, ArbitraryBindingNumbers) {
         })");
 
     wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
-        // TODO(crbug.com/tint/386): Use the same struct
+        [[block]] struct Ubo {
-        [[block]] struct Ubo1 {
             color : vec4<f32>;
         };
 
-        [[block]] struct Ubo2 {
+        [[group(0), binding(953)]] var <uniform> ubo1 : Ubo;
-            color : vec4<f32>;
+        [[group(0), binding(47)]] var <uniform> ubo2 : Ubo;
-        };
+        [[group(0), binding(111)]] var <uniform> ubo3 : Ubo;
-
-        [[block]] struct Ubo3 {
-            color : vec4<f32>;
-        };
-
-        // TODO(crbug.com/tint/386): Use the same struct definition.
-        [[group(0), binding(953)]] var <uniform> ubo1 : Ubo1;
-        [[group(0), binding(47)]] var <uniform> ubo2 : Ubo2;
-        [[group(0), binding(111)]] var <uniform> ubo3 : Ubo3;
 
         [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
             return ubo1.color + 2.0 * ubo2.color + 4.0 * ubo3.color;

src/tests/end2end/ComputeCopyStorageBufferTests.cpp

@@ -88,16 +88,12 @@ void ComputeCopyStorageBufferTests::BasicTest(const char* shader) {
 // Test that a trivial compute-shader memcpy implementation works.
 TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfBasic) {
     BasicTest(R"(
-        [[block]] struct Buf1 {
+        [[block]] struct Buf {
-            s : array<vec4<u32>, 4>;
-        };
-        [[block]] struct Buf2 {
             s : array<vec4<u32>, 4>;
         };
 
-        // TODO(crbug.com/tint/386): Use the same struct type
+        [[group(0), binding(0)]] var<storage, read_write> src : Buf;
-        [[group(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
-        [[group(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute), workgroup_size(1)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -115,16 +111,12 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfStruct) {
             b : vec2<u32>;
         };
 
-        [[block]] struct Buf1 {
+        [[block]] struct Buf {
-            s : array<S, 4>;
-        };
-        [[block]] struct Buf2 {
             s : array<S, 4>;
         };
 
-        // TODO(crbug.com/tint/386): Use the same struct type
+        [[group(0), binding(0)]] var<storage, read_write> src : Buf;
-        [[group(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
-        [[group(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute), workgroup_size(1)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -137,16 +129,12 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfStruct) {
 // Test that a trivial compute-shader memcpy implementation works.
 TEST_P(ComputeCopyStorageBufferTests, UnsizedArrayOfBasic) {
     BasicTest(R"(
-        [[block]] struct Buf1 {
+        [[block]] struct Buf {
-            s : array<vec4<u32>>;
-        };
-        [[block]] struct Buf2 {
             s : array<vec4<u32>>;
         };
 
-        // TODO(crbug.com/tint/386): Use the same struct type
+        [[group(0), binding(0)]] var<storage, read_write> src : Buf;
-        [[group(0), binding(0)]] var<storage, read_write> src : Buf1;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
-        [[group(0), binding(1)]] var<storage, read_write> dst : Buf2;
 
         [[stage(compute), workgroup_size(1)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {

src/tests/end2end/ComputeStorageBufferBarrierTests.cpp

@@ -82,17 +82,12 @@ TEST_P(ComputeStorageBufferBarrierTests, AddPingPong) {
         device, data.data(), bufferSize, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
 
     wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
-        // TODO(crbug.com/tint/386):  Use the same struct.
+        [[block]] struct Buf {
-        [[block]] struct Src {
             data : array<u32, 100>;
         };
 
-        [[block]] struct Dst {
+        [[group(0), binding(0)]] var<storage, read_write> src : Buf;
-            data : array<u32, 100>;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
-        };
-
-        [[group(0), binding(0)]] var<storage, read_write> src : Src;
-        [[group(0), binding(1)]] var<storage, read_write> dst : Dst;
 
         [[stage(compute), workgroup_size(1)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
@@ -153,17 +148,12 @@ TEST_P(ComputeStorageBufferBarrierTests, StorageAndReadonlyStoragePingPongInOneP
         device, data.data(), bufferSize, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
 
     wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
-        // TODO(crbug.com/tint/386):  Use the same struct.
+        [[block]] struct Buf {
-        [[block]] struct Src {
             data : array<u32, 100>;
         };
 
-        [[block]] struct Dst {
+        [[group(0), binding(0)]] var<storage, read> src : Buf;
-            data : array<u32, 100>;
+        [[group(0), binding(1)]] var<storage, read_write> dst : Buf;
-        };
-
-        [[group(0), binding(0)]] var<storage, read> src : Src;
-        [[group(0), binding(1)]] var<storage, read_write> dst : Dst;
 
         [[stage(compute), workgroup_size(1)]]
         fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {

src/tests/end2end/DynamicBufferOffsetTests.cpp

@@ -107,36 +107,19 @@ class DynamicBufferOffsetTests : public DawnTest {
         std::ostringstream fs;
         std::string multipleNumber = isInheritedPipeline ? "2" : "1";
         fs << R"(
-            // TODO(crbug.com/tint/386):  Use the same struct.
+            [[block]] struct Buf {
-            [[block]] struct Buffer1 {
                 value : vec2<u32>;
             };
 
-            [[block]] struct Buffer2 {
+            [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buf;
-                value : vec2<u32>;
+            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buf;
-            };
+            [[group(0), binding(3)]] var<uniform> uBuffer : Buf;
-
+            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buf;
-            [[block]] struct Buffer3 {
-                value : vec2<u32>;
-            };
-
-            [[block]] struct Buffer4 {
-                value : vec2<u32>;
-            };
-
-            [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
-            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2;
-            [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
-            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4;
         )";
 
         if (isInheritedPipeline) {
             fs << R"(
-                [[block]] struct Buffer5 {
+                [[group(1), binding(0)]] var<uniform> paddingBlock : Buf;
-                    value : vec2<u32>;
-                };
-
-                [[group(1), binding(0)]] var<uniform> paddingBlock : Buffer5;
             )";
         }
 
@@ -144,7 +127,7 @@ class DynamicBufferOffsetTests : public DawnTest {
         fs << R"(
             [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> {
                 sBufferNotDynamic.value = uBufferNotDynamic.value.xy;
-                sBuffer.value = vec2<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + sBufferNotDynamic.value.xy);
+                sBuffer.value = vec2<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + uBufferNotDynamic.value.xy);
                 return vec4<f32>(f32(uBuffer.value.x) / 255.0, f32(uBuffer.value.y) / 255.0,
                                       1.0, 1.0);
             }
@@ -174,36 +157,19 @@ class DynamicBufferOffsetTests : public DawnTest {
         std::ostringstream cs;
         std::string multipleNumber = isInheritedPipeline ? "2" : "1";
         cs << R"(
-            // TODO(crbug.com/tint/386):  Use the same struct.
+            [[block]] struct Buf {
-            [[block]] struct Buffer1 {
                 value : vec2<u32>;
             };
 
-            [[block]] struct Buffer2 {
+            [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buf;
-                value : vec2<u32>;
+            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buf;
-            };
+            [[group(0), binding(3)]] var<uniform> uBuffer : Buf;
-
+            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buf;
-            [[block]] struct Buffer3 {
-                value : vec2<u32>;
-            };
-
-            [[block]] struct Buffer4 {
-                value : vec2<u32>;
-            };
-
-            [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
-            [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2;
-            [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
-            [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4;
         )";
 
         if (isInheritedPipeline) {
             cs << R"(
-                [[block]] struct Buffer5 {
+                [[group(1), binding(0)]] var<uniform> paddingBlock : Buf;
-                    value : vec2<u32>;
-                };
-
-                [[group(1), binding(0)]] var<uniform> paddingBlock : Buffer5;
             )";
         }
 
@@ -211,7 +177,7 @@ class DynamicBufferOffsetTests : public DawnTest {
         cs << R"(
             [[stage(compute), workgroup_size(1)]] fn main() {
                 sBufferNotDynamic.value = uBufferNotDynamic.value.xy;
-                sBuffer.value = vec2<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + sBufferNotDynamic.value.xy);
+                sBuffer.value = vec2<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + uBufferNotDynamic.value.xy);
             }
         )";
 
@@ -256,7 +222,7 @@ TEST_P(DynamicBufferOffsetTests, BasicRenderPipeline) {
 }
 
 // Have non-zero dynamic offsets.
-TEST_P(DynamicBufferOffsetTests, SetDynamicOffestsRenderPipeline) {
+TEST_P(DynamicBufferOffsetTests, SetDynamicOffsetsRenderPipeline) {
     wgpu::RenderPipeline pipeline = CreateRenderPipeline();
     utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize);
 
@@ -280,6 +246,9 @@ TEST_P(DynamicBufferOffsetTests, SetDynamicOffestsRenderPipeline) {
 
 // Dynamic offsets are all zero and no effect to result.
 TEST_P(DynamicBufferOffsetTests, BasicComputePipeline) {
+    // TODO(crbug.com/dawn/978): Failing on Windows Vulkan NVIDIA
+    DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsNvidia());
+
     wgpu::ComputePipeline pipeline = CreateComputePipeline();
 
     std::array<uint32_t, 2> offsets = {0, 0};
@@ -298,7 +267,10 @@ TEST_P(DynamicBufferOffsetTests, BasicComputePipeline) {
 }
 
 // Have non-zero dynamic offsets.
-TEST_P(DynamicBufferOffsetTests, SetDynamicOffestsComputePipeline) {
+TEST_P(DynamicBufferOffsetTests, SetDynamicOffsetsComputePipeline) {
+    // TODO(crbug.com/dawn/978): Failing on Windows Vulkan NVIDIA
+    DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsNvidia());
+
     wgpu::ComputePipeline pipeline = CreateComputePipeline();
 
     std::array<uint32_t, 2> offsets = {kMinUniformBufferOffsetAlignment,
@@ -319,7 +291,7 @@ TEST_P(DynamicBufferOffsetTests, SetDynamicOffestsComputePipeline) {
 }
 
 // Test inherit dynamic offsets on render pipeline
-TEST_P(DynamicBufferOffsetTests, InheritDynamicOffestsRenderPipeline) {
+TEST_P(DynamicBufferOffsetTests, InheritDynamicOffsetsRenderPipeline) {
     // Using default pipeline and setting dynamic offsets
     wgpu::RenderPipeline pipeline = CreateRenderPipeline();
     wgpu::RenderPipeline testPipeline = CreateRenderPipeline(true);
@@ -351,7 +323,7 @@ TEST_P(DynamicBufferOffsetTests, InheritDynamicOffestsRenderPipeline) {
 // TODO(shaobo.yan@intel.com) : Try this test on GTX1080 and cannot reproduce the failure.
 // Suspect it is due to dawn doesn't handle sync between two dispatch and disable this case.
 // Will double check root cause after got GTX1660.
-TEST_P(DynamicBufferOffsetTests, InheritDynamicOffestsComputePipeline) {
+TEST_P(DynamicBufferOffsetTests, InheritDynamicOffsetsComputePipeline) {
     DAWN_SUPPRESS_TEST_IF(IsWindows());
     wgpu::ComputePipeline pipeline = CreateComputePipeline();
     wgpu::ComputePipeline testPipeline = CreateComputePipeline(true);
@@ -377,7 +349,7 @@ TEST_P(DynamicBufferOffsetTests, InheritDynamicOffestsComputePipeline) {
 }
 
 // Setting multiple dynamic offsets for the same bindgroup in one render pass.
-TEST_P(DynamicBufferOffsetTests, UpdateDynamicOffestsMultipleTimesRenderPipeline) {
+TEST_P(DynamicBufferOffsetTests, UpdateDynamicOffsetsMultipleTimesRenderPipeline) {
     // Using default pipeline and setting dynamic offsets
     wgpu::RenderPipeline pipeline = CreateRenderPipeline();
 

src/tests/end2end/GpuMemorySynchronizationTests.cpp

@@ -519,15 +519,11 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
         };
         [[group(0), binding(1)]] var<storage, read_write> ibContents : IBContents;
 
-        // TODO(crbug.com/tint/386): Use the same struct.
+        [[block]] struct ColorContents {
-        [[block]] struct ColorContents1 {
             color : f32;
         };
-        [[block]] struct ColorContents2 {
+        [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents;
-            color : f32;
+        [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents;
-        };
-        [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents1;
-        [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents2;
 
         [[stage(compute), workgroup_size(1)]] fn main() {
             vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);

src/tests/end2end/OpArrayLengthTests.cpp

@@ -53,20 +53,15 @@ class OpArrayLengthTest : public DawnTest {
         // Common shader code to use these buffers in shaders, assuming they are in bindgroup index
         // 0.
         mShaderInterface = R"(
-            // TODO(crbug.com/tint/386): Use the same struct.
+            [[block]] struct DataBuffer {
-            [[block]] struct DataBuffer1 {
-                data : [[stride(4)]] array<f32>;
-            };
-
-            [[block]] struct DataBuffer2 {
                 data : [[stride(4)]] array<f32>;
             };
 
             // The length should be 1 because the buffer is 4-byte long.
-            [[group(0), binding(0)]] var<storage, read> buffer1 : DataBuffer1;
+            [[group(0), binding(0)]] var<storage, read> buffer1 : DataBuffer;
 
             // The length should be 64 because the buffer is 256 bytes long.
-            [[group(0), binding(1)]] var<storage, read> buffer2 : DataBuffer2;
+            [[group(0), binding(1)]] var<storage, read> buffer2 : DataBuffer;
 
             // The length should be (512 - 16*4) / 8 = 56 because the buffer is 512 bytes long
             // and the structure is 8 bytes big.