Add end2end tests for compute flow control

This replicates a subset of the flow control CTS tests. The purpose of
this is to provide relatively easy-to-repro cases to demonstrate flow
control bugs on Intel Graphics UHD 630 GPUs.

Bug: tint:1868
Change-Id: I34d692230b44d8a0a917dc773cc748bbf288d55a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/122940
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

src/dawn/tests/BUILD.gn

@@ -493,6 +493,7 @@ source_set("end2end_tests_sources") {
     "end2end/CompressedTextureFormatTests.cpp",
     "end2end/ComputeCopyStorageBufferTests.cpp",
     "end2end/ComputeDispatchTests.cpp",
+    "end2end/ComputeFlowControlTests.cpp",
     "end2end/ComputeLayoutMemoryBufferTests.cpp",
     "end2end/ComputeSharedMemoryTests.cpp",
     "end2end/ComputeStorageBufferBarrierTests.cpp",

src/dawn/tests/end2end/ComputeFlowControlTests.cpp

@@ -0,0 +1,507 @@
+// Copyright 2023 The Dawn Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <vector>
+
+#include "dawn/tests/DawnTest.h"
+
+#include "dawn/utils/WGPUHelpers.h"
+
+// Tests flow control in WGSL shaders. This helps to identify bugs either in Tint's WGSL
+// compilation, or driver shader compilation.
+class ComputeFlowControlTests : public DawnTest {
+  public:
+    void RunTest(const char* shader,
+                 const std::vector<uint32_t>& inputs,
+                 const std::vector<uint32_t>& expected);
+};
+
+void ComputeFlowControlTests::RunTest(const char* shader,
+                                      const std::vector<uint32_t>& inputs,
+                                      const std::vector<uint32_t>& expected) {
+    // Set up shader and pipeline
+    auto module = utils::CreateShaderModule(device, shader);
+
+    wgpu::ComputePipelineDescriptor csDesc;
+    csDesc.compute.module = module;
+    csDesc.compute.entryPoint = "main";
+
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+    // Set up src storage buffer
+    wgpu::Buffer src = utils::CreateBufferFromData(
+        device, inputs.data(), inputs.size() * sizeof(uint32_t),
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst);
+
+    // Set up dst storage buffer
+    std::vector<uint32_t> dst_init_values(expected.size(), 0xDEADBEEF);
+    dst_init_values[0] = 0;  // initial count
+
+    wgpu::Buffer dst = utils::CreateBufferFromData(
+        device, dst_init_values.data(), dst_init_values.size() * sizeof(uint32_t),
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst);
+
+    // Set up bind group and issue dispatch
+    wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+                                                     {
+                                                         {0, src},
+                                                         {1, dst},
+                                                     });
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.DispatchWorkgroups(1);
+        pass.End();
+
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), dst, 0, expected.size());
+}
+
+// Test no branching with one call to push_output
+TEST_P(ComputeFlowControlTests, One) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // ignored
+    };
+    auto expected = std::vector<uint32_t>{1,            // count
+                                          0xA0,         // first
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Test no branching with two calls to push_output
+TEST_P(ComputeFlowControlTests, Two) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  push_output(0xA1);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // ignored
+    };
+    auto expected = std::vector<uint32_t>{2,            // count
+                                          0xA0,         // first
+                                          0xA1,         // second
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Test no branching with three calls to push_output
+TEST_P(ComputeFlowControlTests, Three) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  push_output(0xA1);
+  push_output(0xA2);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // ignored
+    };
+    auto expected = std::vector<uint32_t>{3,            // count
+                                          0xA0,         // first
+                                          0xA1,         // second
+                                          0xA2,         // third
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Test if statement with branch taken
+TEST_P(ComputeFlowControlTests, IfTrue) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        1  // take branch
+    };
+    auto expected = std::vector<uint32_t>{3,            // count
+                                          0xA0,         // before if-else
+                                          0xA1,         // branch
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+
+    RunTest(shader, inputs, expected);
+}
+
+// Test if statement with branch not taken
+TEST_P(ComputeFlowControlTests, IfFalse) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // don't take branch
+    };
+    auto expected = std::vector<uint32_t>{2,            // count
+                                          0xA0,         // before if-else
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+
+    RunTest(shader, inputs, expected);
+}
+
+// Same as IfFalse test, but with push_output calls inlined
+TEST_P(ComputeFlowControlTests, IfFalseInlined) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  {
+    let i = outputs.count;
+    outputs.data[i] = 0xA0u;
+    outputs.count++;
+  }
+
+  if (inputs[0] != 0) {
+      let i = outputs.count;
+      outputs.data[i] = 0xA1u;
+      outputs.count++;
+  }
+
+  {
+    var i = outputs.count;
+    outputs.data[i] = 0xA3u;
+    outputs.count++;
+  }
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // don't take branch
+    };
+    auto expected = std::vector<uint32_t>{2,            // count
+                                          0xA0,         // before if-else
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Same as IfFalse test, but with fixed-size storage arrays
+TEST_P(ComputeFlowControlTests, IfFalseFixedSizeArrays) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32, 2>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32, 1>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // don't take branch
+    };
+    auto expected = std::vector<uint32_t>{2,            // count
+                                          0xA0,         // before if-else
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Same as IfFalse test, but `outputs.count++` is replaced by `outputs.count = i + 1`
+TEST_P(ComputeFlowControlTests, IfFalseNoCountPlusPlus) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count = i + 1;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // don't take branch
+    };
+    auto expected = std::vector<uint32_t>{2,            // count
+                                          0xA0,         // before if-else
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Same as IfFalse test, but `outputs.count++` is replaced by `outputs.count += 4`
+TEST_P(ComputeFlowControlTests, IfFalseIncCountByFour) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count += 4;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // don't take branch
+    };
+    const uint32_t D = 0xDEADBEEF;
+    auto expected = std::vector<uint32_t>{8,               // count
+                                          0xA0, D, D, D,   // before if-else
+                                          0xA3, D, D, D};  // after if-else
+    RunTest(shader, inputs, expected);
+}
+
+// Test if-else statement with true branch taken
+TEST_P(ComputeFlowControlTests, IfElseTrue) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  } else {
+    push_output(0xA2);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        1  // take true branch
+    };
+    auto expected = std::vector<uint32_t>{3,            // count
+                                          0xA0,         // before if-else
+                                          0xA1,         // true branch
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+// Test if-else statement with false branch taken
+TEST_P(ComputeFlowControlTests, IfElseFalse) {
+    const char* shader = R"(
+struct Outputs {
+  count : u32,
+  data  : array<u32>,
+};
+@group(0) @binding(0) var<storage, read>       inputs  : array<u32>;
+@group(0) @binding(1) var<storage, read_write> outputs : Outputs;
+
+fn push_output(value : u32) {
+  let i = outputs.count;
+  outputs.data[i] = value;
+  outputs.count++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  _ = &inputs;
+  _ = &outputs;
+
+  push_output(0xA0);
+  if (inputs[0] != 0) {
+    push_output(0xA1);
+  } else {
+    push_output(0xA2);
+  }
+  push_output(0xA3);
+})";
+
+    auto inputs = std::vector<uint32_t>{
+        0  // take false branch
+    };
+    auto expected = std::vector<uint32_t>{3,            // count
+                                          0xA0,         // before if-else
+                                          0xA2,         // false branch
+                                          0xA3,         // after if-else
+                                          0xDEADBEEF};  // unwritten
+    RunTest(shader, inputs, expected);
+}
+
+DAWN_INSTANTIATE_TEST(ComputeFlowControlTests,
+                      D3D12Backend(),
+                      MetalBackend(),
+                      OpenGLBackend(),
+                      OpenGLESBackend(),
+                      VulkanBackend());