Add tests for some limits

tested:
 - maxComputeWorkgroupStorageSize
 - maxUniformBufferBindingSize
 - maxStorageBufferBindingSize

Two of these limits are exposed as configurable to the JS API
so it's important they are tested to work before we expose
them. maxUniformBufferBindingSize came along as well because
the test for storageBufferBindingSize was easy to parameterize.

Bug: dawn:685
Change-Id: I08de6df9d70a22aca0f48ac3fef0038f7aec727b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/66480
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Brandon Jones <bajones@chromium.org>
Commit-Queue: Austin Eng <enga@chromium.org>

src/tests/BUILD.gn

@@ -355,6 +355,7 @@ source_set("dawn_end2end_tests_sources") {
     "end2end/FirstIndexOffsetTests.cpp",
     "end2end/GpuMemorySynchronizationTests.cpp",
     "end2end/IndexFormatTests.cpp",
+    "end2end/MaxLimitTests.cpp",
     "end2end/MemoryAllocationStressTests.cpp",
     "end2end/MultisampledRenderingTests.cpp",
     "end2end/MultisampledSamplingTests.cpp",

src/tests/DawnTest.cpp

@@ -858,6 +858,10 @@ std::vector<const char*> DawnTestBase::GetRequiredFeatures() {
     return {};
 }
 
+wgpu::RequiredLimits DawnTestBase::GetRequiredLimits(const wgpu::SupportedLimits&) {
+    return {};
+}
+
 const wgpu::AdapterProperties& DawnTestBase::GetAdapterProperties() const {
     return mParam.adapterProperties;
 }
@@ -921,6 +925,11 @@ void DawnTestBase::SetUp() {
     deviceDescriptor.forceDisabledToggles = mParam.forceDisabledWorkarounds;
     deviceDescriptor.requiredFeatures = GetRequiredFeatures();
 
+    wgpu::SupportedLimits supportedLimits;
+    mBackendAdapter.GetLimits(reinterpret_cast<WGPUSupportedLimits*>(&supportedLimits));
+    wgpu::RequiredLimits requiredLimits = GetRequiredLimits(supportedLimits);
+    deviceDescriptor.requiredLimits = reinterpret_cast<WGPURequiredLimits*>(&requiredLimits);
+
     // Disabled disallowing unsafe APIs so we can test them.
     deviceDescriptor.forceDisabledToggles.push_back("disallow_unsafe_apis");
 

src/tests/DawnTest.h

@@ -483,6 +483,8 @@ class DawnTestBase {
     // code path to handle the situation when not all features are supported.
     virtual std::vector<const char*> GetRequiredFeatures();
 
+    virtual wgpu::RequiredLimits GetRequiredLimits(const wgpu::SupportedLimits&);
+
     const wgpu::AdapterProperties& GetAdapterProperties() const;
 
     // TODO(crbug.com/dawn/689): Use limits returned from the wire

src/tests/end2end/MaxLimitTests.cpp

@@ -0,0 +1,226 @@
+// Copyright 2021 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 "tests/DawnTest.h"
+
+#include "common/Math.h"
+#include "utils/WGPUHelpers.h"
+
+class MaxLimitTests : public DawnTest {
+  public:
+    wgpu::RequiredLimits GetRequiredLimits(const wgpu::SupportedLimits& supported) override {
+        wgpu::RequiredLimits required = {};
+        required.limits = supported.limits;
+        return required;
+    }
+};
+
+// Test using the maximum amount of workgroup memory works
+TEST_P(MaxLimitTests, MaxComputeWorkgroupStorageSize) {
+    uint32_t maxComputeWorkgroupStorageSize =
+        GetSupportedLimits().limits.maxComputeWorkgroupStorageSize;
+
+    std::string shader = R"(
+        [[block]] struct Dst {
+            value0 : u32;
+            value1 : u32;
+        };
+
+        [[group(0), binding(0)]] var<storage, write> dst : Dst;
+
+        struct WGData {
+          value0 : u32;
+          // padding such that value0 and value1 are the first and last bytes of the memory.
+          [[size()" + std::to_string(maxComputeWorkgroupStorageSize / 4 - 2) +
+                         R"()]] padding : u32;
+          value1 : u32;
+        };
+        var<workgroup> wg_data : WGData;
+
+        [[stage(compute), workgroup_size(2,1,1)]]
+        fn main([[builtin(local_invocation_index)]] LocalInvocationIndex : u32) {
+            if (LocalInvocationIndex == 0u) {
+                // Put data into the first and last byte of workgroup memory.
+                wg_data.value0 = 79u;
+                wg_data.value1 = 42u;
+            }
+
+            workgroupBarrier();
+
+            if (LocalInvocationIndex == 1u) {
+                // Read data out of workgroup memory into a storage buffer.
+                dst.value0 = wg_data.value0;
+                dst.value1 = wg_data.value1;
+            }
+        }
+    )";
+    wgpu::ComputePipelineDescriptor csDesc;
+    csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
+    csDesc.compute.entryPoint = "main";
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+    // Set up dst storage buffer
+    wgpu::BufferDescriptor dstDesc;
+    dstDesc.size = 8;
+    dstDesc.usage =
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
+    wgpu::Buffer dst = device.CreateBuffer(&dstDesc);
+
+    // Set up bind group and issue dispatch
+    wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+                                                     {
+                                                         {0, dst},
+                                                     });
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bindGroup);
+    pass.Dispatch(1);
+    pass.EndPass();
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_EQ(79, dst, 0);
+    EXPECT_BUFFER_U32_EQ(42, dst, 4);
+}
+
+// Test using the maximum uniform/storage buffer binding size works
+TEST_P(MaxLimitTests, MaxBufferBindingSize) {
+    // The uniform buffer layout used in this test is not supported on ES.
+    DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
+
+    for (wgpu::BufferUsage usage : {wgpu::BufferUsage::Storage, wgpu::BufferUsage::Uniform}) {
+        uint64_t maxBufferBindingSize;
+        std::string shader;
+        switch (usage) {
+            case wgpu::BufferUsage::Storage:
+                maxBufferBindingSize = GetSupportedLimits().limits.maxStorageBufferBindingSize;
+                // TODO(crbug.com/dawn/1160): Usually can't actually allocate a buffer this large
+                // because allocating the buffer for zero-initialization fails.
+                maxBufferBindingSize =
+                    std::min(maxBufferBindingSize, uint64_t(2) * 1024 * 1024 * 1024);
+                if (IsWARP()) {
+                    maxBufferBindingSize =
+                        std::min(maxBufferBindingSize, uint64_t(1) * 1024 * 1024 * 1024);
+                }
+                shader = R"(
+                  [[block]] struct Buf {
+                      value0 : u32;
+                      // padding such that value0 and value1 are the first and last bytes of the memory.
+                      [[size()" +
+                         std::to_string(maxBufferBindingSize - 8) + R"()]] padding : u32;
+                      value1 : u32;
+                  };
+
+                  [[block]] struct Result {
+                      value0 : u32;
+                      value1 : u32;
+                  };
+
+                  [[group(0), binding(0)]] var<storage, read> buf : Buf;
+                  [[group(0), binding(1)]] var<storage, write> result : Result;
+
+                  [[stage(compute), workgroup_size(1,1,1)]]
+                  fn main() {
+                      result.value0 = buf.value0;
+                      result.value1 = buf.value1;
+                  }
+              )";
+                break;
+            case wgpu::BufferUsage::Uniform:
+                maxBufferBindingSize = GetSupportedLimits().limits.maxUniformBufferBindingSize;
+                shader = R"(
+                  [[block]] struct Buf {
+                      value0 : u32;
+                      // padding such that value0 and value1 are the first and last bytes of the memory.
+                      [[size()" +
+                         std::to_string(maxBufferBindingSize - 8) + R"()]] padding : u32;
+                      value1 : u32;
+                  };
+
+                  [[block]] struct Result {
+                      value0 : u32;
+                      value1 : u32;
+                  };
+
+                  [[group(0), binding(0)]] var<uniform> buf : Buf;
+                  [[group(0), binding(1)]] var<storage, write> result : Result;
+
+                  [[stage(compute), workgroup_size(1,1,1)]]
+                  fn main() {
+                      result.value0 = buf.value0;
+                      result.value1 = buf.value1;
+                  }
+              )";
+                break;
+            default:
+                UNREACHABLE();
+        }
+
+        device.PushErrorScope(wgpu::ErrorFilter::OutOfMemory);
+
+        wgpu::BufferDescriptor bufDesc;
+        bufDesc.size = maxBufferBindingSize;
+        bufDesc.usage = usage | wgpu::BufferUsage::CopyDst;
+        wgpu::Buffer buffer = device.CreateBuffer(&bufDesc);
+
+        WGPUErrorType oomResult;
+        device.PopErrorScope([](WGPUErrorType type, const char*,
+                                void* userdata) { *static_cast<WGPUErrorType*>(userdata) = type; },
+                             &oomResult);
+        FlushWire();
+        // Max buffer size is smaller than the max buffer binding size.
+        DAWN_TEST_UNSUPPORTED_IF(oomResult == WGPUErrorType_OutOfMemory);
+
+        wgpu::BufferDescriptor resultBufDesc;
+        resultBufDesc.size = 8;
+        resultBufDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+        wgpu::Buffer resultBuffer = device.CreateBuffer(&resultBufDesc);
+
+        uint32_t value0 = 89234;
+        queue.WriteBuffer(buffer, 0, &value0, sizeof(value0));
+
+        uint32_t value1 = 234;
+        uint64_t value1Offset = Align(maxBufferBindingSize - sizeof(value1), 4);
+        queue.WriteBuffer(buffer, value1Offset, &value1, sizeof(value1));
+
+        wgpu::ComputePipelineDescriptor csDesc;
+        csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
+        csDesc.compute.entryPoint = "main";
+        wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+        wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+                                                         {{0, buffer}, {1, resultBuffer}});
+
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.Dispatch(1);
+        pass.EndPass();
+        wgpu::CommandBuffer commands = encoder.Finish();
+        queue.Submit(1, &commands);
+
+        EXPECT_BUFFER_U32_EQ(value0, resultBuffer, 0);
+        EXPECT_BUFFER_U32_EQ(value1, resultBuffer, 4);
+    }
+}
+
+DAWN_INSTANTIATE_TEST(MaxLimitTests,
+                      D3D12Backend(),
+                      MetalBackend(),
+                      OpenGLBackend(),
+                      OpenGLESBackend(),
+                      VulkanBackend());