diff --git a/src/tests/BUILD.gn b/src/tests/BUILD.gn index 3f236fe9c0..e1e658a590 100644 --- a/src/tests/BUILD.gn +++ b/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", diff --git a/src/tests/DawnTest.cpp b/src/tests/DawnTest.cpp index 0dff12615e..cb597751c2 100644 --- a/src/tests/DawnTest.cpp +++ b/src/tests/DawnTest.cpp @@ -858,6 +858,10 @@ std::vector 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(&supportedLimits)); + wgpu::RequiredLimits requiredLimits = GetRequiredLimits(supportedLimits); + deviceDescriptor.requiredLimits = reinterpret_cast(&requiredLimits); + // Disabled disallowing unsafe APIs so we can test them. deviceDescriptor.forceDisabledToggles.push_back("disallow_unsafe_apis"); diff --git a/src/tests/DawnTest.h b/src/tests/DawnTest.h index 2d96e1f9b7..db13ff64e3 100644 --- a/src/tests/DawnTest.h +++ b/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 GetRequiredFeatures(); + virtual wgpu::RequiredLimits GetRequiredLimits(const wgpu::SupportedLimits&); + const wgpu::AdapterProperties& GetAdapterProperties() const; // TODO(crbug.com/dawn/689): Use limits returned from the wire diff --git a/src/tests/end2end/MaxLimitTests.cpp b/src/tests/end2end/MaxLimitTests.cpp new file mode 100644 index 0000000000..a083cf0862 --- /dev/null +++ b/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 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 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 buf : Buf; + [[group(0), binding(1)]] var 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 buf : Buf; + [[group(0), binding(1)]] var 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(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());