| // Copyright 2019 The Dawn & Tint Authors | |
| // | |
| // Redistribution and use in source and binary forms, with or without | |
| // modification, are permitted provided that the following conditions are met: | |
| // | |
| // 1. Redistributions of source code must retain the above copyright notice, this | |
| // list of conditions and the following disclaimer. | |
| // | |
| // 2. Redistributions in binary form must reproduce the above copyright notice, | |
| // this list of conditions and the following disclaimer in the documentation | |
| // and/or other materials provided with the distribution. | |
| // | |
| // 3. Neither the name of the copyright holder nor the names of its | |
| // contributors may be used to endorse or promote products derived from | |
| // this software without specific prior written permission. | |
| // | |
| // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" | |
| // AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE | |
| // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE | |
| // DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE | |
| // FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL | |
| // DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR | |
| // SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER | |
| // CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, | |
| // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE | |
| // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. | |
| #include <algorithm> | |
| #include <numeric> | |
| #include <string> | |
| #include <vector> | |
| #include "src/dawn/common/Math.h" | |
| #include "src/dawn/tests/DawnTest.h" | |
| #include "src/dawn/utils/ComboRenderPipelineDescriptor.h" | |
| #include "src/dawn/utils/WGPUHelpers.h" | |
| #include "src/utils/compiler.h" | |
| namespace dawn { | |
| namespace { | |
| constexpr uint32_t kRTSize = 400; | |
| constexpr uint32_t kBindingSize = 8; | |
| class DynamicBufferOffsetTests : public DawnTest { | |
| protected: | |
| uint32_t mMinUniformBufferOffsetAlignment; | |
| uint32_t mMinStorageBufferOffsetAlignment; | |
| wgpu::BindGroup mBindGroups[2]; | |
| wgpu::BindGroupLayout mBindGroupLayouts[2]; | |
| // All buffers are usually 256 bytes / 4 + 2 = 66 u32 elements | |
| // | |
| // mUniformBuffers[0]: | |
| // - not dynamic | |
| // - @group(0) @binding(0) | |
| // - [1, 2, 0 ... 0] | |
| // mUniformBuffers[1]: | |
| // - dynamic | |
| // - @group(0) @binding(3) | |
| // - [1, 2, 0, ..., 0, 5, 6] | |
| // mUniformBuffers[2]: | |
| // - not dynamic | |
| // - @group(1) @binding(0) | |
| // - [1, 2, 0 ... 0] (used for inheritance tests) | |
| wgpu::Buffer mUniformBuffers[3]; | |
| // mStorageBuffers[0]: | |
| // - not dynamic | |
| // - @group(0) @binding(1) | |
| // - shader sets this.xy to mUniformBuffers[0].xy, so [1, 2] | |
| // mStorageBuffers[1]: | |
| // - dynamic | |
| // - @group(0) @binding(4) | |
| // - shader sets this.xy to (mUniformBuffers[0].xy + mUniformBuffers[1]) * multipleNumber | |
| wgpu::Buffer mStorageBuffers[2]; | |
| wgpu::Texture mColorAttachment; | |
| void GetRequiredLimits(const dawn::utils::ComboLimits& supported, | |
| dawn::utils::ComboLimits& required) override { | |
| // TODO(crbug.com/383593270): Enable all the limits. | |
| required.maxStorageBuffersInFragmentStage = supported.maxStorageBuffersInFragmentStage; | |
| required.maxStorageBuffersPerShaderStage = supported.maxStorageBuffersPerShaderStage; | |
| required.minStorageBufferOffsetAlignment = supported.minStorageBufferOffsetAlignment; | |
| } | |
| void SetUp() override { | |
| DawnTest::SetUp(); | |
| DAWN_TEST_UNSUPPORTED_IF(GetSupportedLimits().maxStorageBuffersInFragmentStage < 1); | |
| mMinUniformBufferOffsetAlignment = GetSupportedLimits().minUniformBufferOffsetAlignment; | |
| mMinStorageBufferOffsetAlignment = GetSupportedLimits().minStorageBufferOffsetAlignment; | |
| // Mix up dynamic and non dynamic resources in one bind group and using not continuous | |
| // binding number to cover more cases. | |
| const uint32_t maxAlignment = | |
| std::max(mMinUniformBufferOffsetAlignment, mMinStorageBufferOffsetAlignment); | |
| std::vector<uint32_t> uniformData(maxAlignment / sizeof(uint32_t) + 2); | |
| uniformData[0] = 1; | |
| uniformData[1] = 2; | |
| mUniformBuffers[0] = utils::CreateBufferFromData(device, uniformData.data(), | |
| sizeof(uint32_t) * uniformData.size(), | |
| wgpu::BufferUsage::Uniform); | |
| uniformData[uniformData.size() - 2] = 5; | |
| uniformData[uniformData.size() - 1] = 6; | |
| // Dynamic uniform buffer | |
| mUniformBuffers[1] = utils::CreateBufferFromData(device, uniformData.data(), | |
| sizeof(uint32_t) * uniformData.size(), | |
| wgpu::BufferUsage::Uniform); | |
| wgpu::BufferDescriptor storageBufferDescriptor; | |
| storageBufferDescriptor.size = sizeof(uint32_t) * uniformData.size(); | |
| storageBufferDescriptor.usage = | |
| wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; | |
| mStorageBuffers[0] = device.CreateBuffer(&storageBufferDescriptor); | |
| // Dynamic storage buffer | |
| mStorageBuffers[1] = device.CreateBuffer(&storageBufferDescriptor); | |
| // Default bind group layout | |
| mBindGroupLayouts[0] = utils::MakeBindGroupLayout( | |
| device, {{0, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Uniform}, | |
| {1, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Storage}, | |
| {3, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Uniform, true}, | |
| {4, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Storage, true}}); | |
| // Default bind group | |
| mBindGroups[0] = utils::MakeBindGroup(device, mBindGroupLayouts[0], | |
| {{0, mUniformBuffers[0], 0, kBindingSize}, | |
| {1, mStorageBuffers[0], 0, kBindingSize}, | |
| {3, mUniformBuffers[1], 0, kBindingSize}, | |
| {4, mStorageBuffers[1], 0, kBindingSize}}); | |
| // Extra uniform buffer for inheriting test | |
| mUniformBuffers[2] = utils::CreateBufferFromData(device, uniformData.data(), | |
| sizeof(uint32_t) * uniformData.size(), | |
| wgpu::BufferUsage::Uniform); | |
| // Bind group layout for inheriting test | |
| mBindGroupLayouts[1] = utils::MakeBindGroupLayout( | |
| device, {{0, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Uniform}}); | |
| // Bind group for inheriting test | |
| mBindGroups[1] = utils::MakeBindGroup(device, mBindGroupLayouts[1], | |
| {{0, mUniformBuffers[2], 0, kBindingSize}}); | |
| } | |
| wgpu::RenderPipeline CreateRenderPipeline(bool isInheritedPipeline = false) { | |
| wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( | |
| @vertex | |
| fn main(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f { | |
| var pos = array( | |
| vec2f(-1.0, 0.0), | |
| vec2f(-1.0, 1.0), | |
| vec2f( 0.0, 1.0)); | |
| return vec4f(pos[VertexIndex], 0.0, 1.0); | |
| })"); | |
| // Construct fragment shader source | |
| std::ostringstream fs; | |
| std::string multipleNumber = isInheritedPipeline ? "2" : "1"; | |
| fs << R"( | |
| struct Buf { | |
| value : vec2u | |
| } | |
| @group(0) @binding(0) var<uniform> uBufferNotDynamic : Buf; | |
| @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; | |
| )"; | |
| if (isInheritedPipeline) { | |
| fs << R"( | |
| @group(1) @binding(0) var<uniform> paddingBlock : Buf; | |
| )"; | |
| } | |
| fs << "const multipleNumber : u32 = " << multipleNumber << "u;\n"; | |
| fs << R"( | |
| @fragment fn main() -> @location(0) vec4f { | |
| sBufferNotDynamic.value = uBufferNotDynamic.value.xy; | |
| sBuffer.value = vec2u(multipleNumber, multipleNumber) * (uBuffer.value.xy + uBufferNotDynamic.value.xy); | |
| return vec4f(f32(uBuffer.value.x) / 255.0, f32(uBuffer.value.y) / 255.0, | |
| 1.0, 1.0); | |
| } | |
| )"; | |
| wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fs.str().c_str()); | |
| utils::ComboRenderPipelineDescriptor pipelineDescriptor; | |
| pipelineDescriptor.vertex.module = vsModule; | |
| pipelineDescriptor.cFragment.module = fsModule; | |
| pipelineDescriptor.cTargets[0].format = wgpu::TextureFormat::RGBA8Unorm; | |
| wgpu::PipelineLayoutDescriptor pipelineLayoutDescriptor; | |
| if (isInheritedPipeline) { | |
| pipelineLayoutDescriptor.bindGroupLayoutCount = 2; | |
| } else { | |
| pipelineLayoutDescriptor.bindGroupLayoutCount = 1; | |
| } | |
| pipelineLayoutDescriptor.bindGroupLayouts = mBindGroupLayouts; | |
| pipelineDescriptor.layout = device.CreatePipelineLayout(&pipelineLayoutDescriptor); | |
| return device.CreateRenderPipeline(&pipelineDescriptor); | |
| } | |
| wgpu::ComputePipeline CreateComputePipeline(bool isInheritedPipeline = false) { | |
| // Construct compute shader source | |
| std::ostringstream cs; | |
| std::string multipleNumber = isInheritedPipeline ? "2" : "1"; | |
| cs << R"( | |
| struct Buf { | |
| value : vec2u | |
| } | |
| @group(0) @binding(0) var<uniform> uBufferNotDynamic : Buf; | |
| @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; | |
| )"; | |
| if (isInheritedPipeline) { | |
| cs << R"( | |
| @group(1) @binding(0) var<uniform> paddingBlock : Buf; | |
| )"; | |
| } | |
| cs << "const multipleNumber : u32 = " << multipleNumber << "u;\n"; | |
| cs << R"( | |
| @compute @workgroup_size(1) fn main() { | |
| sBufferNotDynamic.value = uBufferNotDynamic.value.xy; | |
| sBuffer.value = vec2u(multipleNumber, multipleNumber) * (uBuffer.value.xy + uBufferNotDynamic.value.xy); | |
| } | |
| )"; | |
| wgpu::ShaderModule csModule = utils::CreateShaderModule(device, cs.str().c_str()); | |
| wgpu::ComputePipelineDescriptor csDesc; | |
| csDesc.compute.module = csModule; | |
| wgpu::PipelineLayoutDescriptor pipelineLayoutDescriptor; | |
| if (isInheritedPipeline) { | |
| pipelineLayoutDescriptor.bindGroupLayoutCount = 2; | |
| } else { | |
| pipelineLayoutDescriptor.bindGroupLayoutCount = 1; | |
| } | |
| pipelineLayoutDescriptor.bindGroupLayouts = mBindGroupLayouts; | |
| csDesc.layout = device.CreatePipelineLayout(&pipelineLayoutDescriptor); | |
| return device.CreateComputePipeline(&csDesc); | |
| } | |
| }; | |
| // Dynamic offsets are all zero and no effect to result. | |
| TEST_P(DynamicBufferOffsetTests, BasicRenderPipeline) { | |
| // TODO(crbug.com/522869941): Produces incorrect result on Pixel 10. | |
| DAWN_SUPPRESS_TEST_IF(IsAndroid() && IsImgTec() && IsVulkan()); | |
| wgpu::RenderPipeline pipeline = CreateRenderPipeline(); | |
| utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| std::array<uint32_t, 2> offsets = {0, 0}; | |
| wgpu::RenderPassEncoder renderPassEncoder = | |
| commandEncoder.BeginRenderPass(&renderPass.renderPassInfo); | |
| renderPassEncoder.SetPipeline(pipeline); | |
| renderPassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {2, 4}; | |
| EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(1, 2, 255, 255), renderPass.color, 0, 0); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], 0, expectedData.size()); | |
| } | |
| // Have non-zero dynamic offsets. | |
| TEST_P(DynamicBufferOffsetTests, SetDynamicOffsetsRenderPipeline) { | |
| // TODO(42242119): fail on Qualcomm Adreno X1. | |
| DAWN_SUPPRESS_TEST_IF(IsD3D11() && IsQualcomm()); | |
| // TODO(crbug.com/522869941): Produces incorrect result on Pixel 10. | |
| DAWN_SUPPRESS_TEST_IF(IsAndroid() && IsImgTec() && IsVulkan()); | |
| wgpu::RenderPipeline pipeline = CreateRenderPipeline(); | |
| utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| wgpu::RenderPassEncoder renderPassEncoder = | |
| commandEncoder.BeginRenderPass(&renderPass.renderPassInfo); | |
| renderPassEncoder.SetPipeline(pipeline); | |
| renderPassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {6, 8}; | |
| EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(5, 6, 255, 255), renderPass.color, 0, 0); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], | |
| mMinStorageBufferOffsetAlignment, expectedData.size()); | |
| } | |
| // Dynamic offsets are all zero and no effect to result. | |
| TEST_P(DynamicBufferOffsetTests, BasicComputePipeline) { | |
| wgpu::ComputePipeline pipeline = CreateComputePipeline(); | |
| std::array<uint32_t, 2> offsets = {0, 0}; | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); | |
| computePassEncoder.SetPipeline(pipeline); | |
| computePassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {2, 4}; | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], 0, expectedData.size()); | |
| } | |
| // Have non-zero dynamic offsets. | |
| TEST_P(DynamicBufferOffsetTests, SetDynamicOffsetsComputePipeline) { | |
| // TODO(42242119): fail on Qualcomm Adreno X1. | |
| DAWN_SUPPRESS_TEST_IF(IsD3D11() && IsQualcomm()); | |
| wgpu::ComputePipeline pipeline = CreateComputePipeline(); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); | |
| computePassEncoder.SetPipeline(pipeline); | |
| computePassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {6, 8}; | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], | |
| mMinStorageBufferOffsetAlignment, expectedData.size()); | |
| } | |
| // Test basic inherit on render pipeline | |
| TEST_P(DynamicBufferOffsetTests, BasicInheritRenderPipeline) { | |
| // TODO(42242119): fail on Qualcomm Adreno X1. | |
| DAWN_SUPPRESS_TEST_IF(IsD3D11() && IsQualcomm()); | |
| // TODO(crbug.com/40238674): Fails on Pixel 10. | |
| DAWN_SUPPRESS_TEST_IF(IsImgTec()); | |
| wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( | |
| @vertex | |
| fn main(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f { | |
| var pos = array( | |
| vec2f(-1.0, 0.0), | |
| vec2f(-1.0, 1.0), | |
| vec2f( 0.0, 1.0)); | |
| return vec4f(pos[VertexIndex], 0.0, 1.0); | |
| })"); | |
| // Construct fragment shader source | |
| std::ostringstream fs; | |
| fs << R"( | |
| struct Buf { | |
| value : vec2u | |
| } | |
| @group(0) @binding(0) var<uniform> uBufferNotDynamic : Buf; | |
| @group(0) @binding(1) var<storage, read_write> sBufferNotDynamic : Buf; | |
| @group(1) @binding(3) var<uniform> uBuffer : Buf; | |
| @group(1) @binding(4) var<storage, read_write> sBuffer : Buf; | |
| @fragment fn main() -> @location(0) vec4f { | |
| sBufferNotDynamic.value = uBufferNotDynamic.value.xy; | |
| sBuffer.value = uBuffer.value.xy + uBufferNotDynamic.value.xy; | |
| return vec4f(f32(uBuffer.value.x) / 255.0, f32(uBuffer.value.y) / 255.0, | |
| 1.0, 1.0); | |
| } | |
| )"; | |
| wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fs.str().c_str()); | |
| utils::ComboRenderPipelineDescriptor pipelineDescriptor; | |
| pipelineDescriptor.vertex.module = vsModule; | |
| pipelineDescriptor.cFragment.module = fsModule; | |
| pipelineDescriptor.cTargets[0].format = wgpu::TextureFormat::RGBA8Unorm; | |
| wgpu::BindGroupLayout bindGroupLayouts[2]; | |
| bindGroupLayouts[0] = utils::MakeBindGroupLayout( | |
| device, {{0, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Uniform}, | |
| {1, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Storage}}); | |
| bindGroupLayouts[1] = utils::MakeBindGroupLayout( | |
| device, {{3, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Uniform, true}, | |
| {4, wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment, | |
| wgpu::BufferBindingType::Storage, true}}); | |
| wgpu::BindGroup bindGroups[2]; | |
| bindGroups[0] = utils::MakeBindGroup(device, bindGroupLayouts[0], | |
| { | |
| {0, mUniformBuffers[0], 0, kBindingSize}, | |
| {1, mStorageBuffers[0], 0, kBindingSize}, | |
| }); | |
| bindGroups[1] = utils::MakeBindGroup( | |
| device, bindGroupLayouts[1], | |
| {{3, mUniformBuffers[1], 0, kBindingSize}, {4, mStorageBuffers[1], 0, kBindingSize}}); | |
| wgpu::PipelineLayoutDescriptor pipelineLayoutDescriptor; | |
| pipelineLayoutDescriptor.bindGroupLayoutCount = 2; | |
| pipelineLayoutDescriptor.bindGroupLayouts = bindGroupLayouts; | |
| pipelineDescriptor.layout = device.CreatePipelineLayout(&pipelineLayoutDescriptor); | |
| wgpu::RenderPipeline pipeline0 = device.CreateRenderPipeline(&pipelineDescriptor); | |
| wgpu::RenderPipeline pipeline1 = device.CreateRenderPipeline(&pipelineDescriptor); | |
| utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| std::array<uint32_t, 2> offsets0 = {0, 0}; | |
| std::array<uint32_t, 2> offsets1 = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| wgpu::RenderPassEncoder renderPassEncoder = | |
| commandEncoder.BeginRenderPass(&renderPass.renderPassInfo); | |
| renderPassEncoder.SetPipeline(pipeline0); | |
| renderPassEncoder.SetBindGroup(0, bindGroups[0]); | |
| renderPassEncoder.SetBindGroup(1, bindGroups[1], offsets0.size(), offsets0.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.SetPipeline(pipeline1); | |
| // bind group 0 should be inherited and still available. | |
| renderPassEncoder.SetBindGroup(1, bindGroups[1], offsets1.size(), offsets1.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {6, 8}; | |
| EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(5, 6, 255, 255), renderPass.color, 0, 0); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], | |
| mMinStorageBufferOffsetAlignment, expectedData.size()); | |
| } | |
| // Test inherit dynamic offsets on render pipeline | |
| TEST_P(DynamicBufferOffsetTests, InheritDynamicOffsetsRenderPipeline) { | |
| // TODO(crbug.com/1497726): Remove when test is no longer flaky on M2 | |
| // devices. | |
| DAWN_SUPPRESS_TEST_IF(IsApple()); | |
| // TODO(crbug.com/40287156): Remove when test is no longer flaky on Pixel 6 | |
| DAWN_SUPPRESS_TEST_IF(IsAndroid() && IsARM()); | |
| // TODO(42242119): fail on Qualcomm Adreno X1. | |
| DAWN_SUPPRESS_TEST_IF(IsD3D11() && IsQualcomm()); | |
| // TODO(crbug.com/522869941): Produces incorrect result on Pixel 10. | |
| DAWN_SUPPRESS_TEST_IF(IsAndroid() && IsImgTec() && IsVulkan()); | |
| // Using default pipeline and setting dynamic offsets | |
| wgpu::RenderPipeline pipeline = CreateRenderPipeline(); | |
| wgpu::RenderPipeline testPipeline = CreateRenderPipeline(true); | |
| utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| wgpu::RenderPassEncoder renderPassEncoder = | |
| commandEncoder.BeginRenderPass(&renderPass.renderPassInfo); | |
| renderPassEncoder.SetPipeline(pipeline); | |
| renderPassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.SetPipeline(testPipeline); | |
| renderPassEncoder.SetBindGroup(1, mBindGroups[1]); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {12, 16}; | |
| EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(5, 6, 255, 255), renderPass.color, 0, 0); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], | |
| mMinStorageBufferOffsetAlignment, expectedData.size()); | |
| } | |
| // Test inherit dynamic offsets on compute pipeline | |
| TEST_P(DynamicBufferOffsetTests, InheritDynamicOffsetsComputePipeline) { | |
| // TODO(42242119): fail on Qualcomm Adreno X1. | |
| DAWN_SUPPRESS_TEST_IF(IsD3D11() && IsQualcomm()); | |
| wgpu::ComputePipeline pipeline = CreateComputePipeline(); | |
| wgpu::ComputePipeline testPipeline = CreateComputePipeline(true); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); | |
| computePassEncoder.SetPipeline(pipeline); | |
| computePassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.SetPipeline(testPipeline); | |
| computePassEncoder.SetBindGroup(1, mBindGroups[1]); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {12, 16}; | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], | |
| mMinStorageBufferOffsetAlignment, expectedData.size()); | |
| } | |
| // Setting multiple dynamic offsets for the same bindgroup in one render pass. | |
| TEST_P(DynamicBufferOffsetTests, UpdateDynamicOffsetsMultipleTimesRenderPipeline) { | |
| // TODO(crbug.com/522869941): Produces incorrect result on Pixel 10. | |
| DAWN_SUPPRESS_TEST_IF(IsAndroid() && IsImgTec() && IsVulkan()); | |
| // Using default pipeline and setting dynamic offsets | |
| wgpu::RenderPipeline pipeline = CreateRenderPipeline(); | |
| utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, kRTSize, kRTSize); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| std::array<uint32_t, 2> testOffsets = {0, 0}; | |
| wgpu::RenderPassEncoder renderPassEncoder = | |
| commandEncoder.BeginRenderPass(&renderPass.renderPassInfo); | |
| renderPassEncoder.SetPipeline(pipeline); | |
| renderPassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.SetBindGroup(0, mBindGroups[0], testOffsets.size(), testOffsets.data()); | |
| renderPassEncoder.Draw(3); | |
| renderPassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {2, 4}; | |
| EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(1, 2, 255, 255), renderPass.color, 0, 0); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], 0, expectedData.size()); | |
| } | |
| // Setting multiple dynamic offsets for the same bindgroup in one compute pass. | |
| TEST_P(DynamicBufferOffsetTests, UpdateDynamicOffsetsMultipleTimesComputePipeline) { | |
| wgpu::ComputePipeline pipeline = CreateComputePipeline(); | |
| std::array<uint32_t, 2> offsets = {mMinUniformBufferOffsetAlignment, | |
| mMinStorageBufferOffsetAlignment}; | |
| std::array<uint32_t, 2> testOffsets = {0, 0}; | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); | |
| computePassEncoder.SetPipeline(pipeline); | |
| computePassEncoder.SetBindGroup(0, mBindGroups[0], offsets.size(), offsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.SetBindGroup(0, mBindGroups[0], testOffsets.size(), testOffsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| std::vector<uint32_t> expectedData = {2, 4}; | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), mStorageBuffers[1], 0, expectedData.size()); | |
| } | |
| namespace { | |
| using ReadBufferUsage = wgpu::BufferUsage; | |
| using OOBRead = bool; | |
| using OOBWrite = bool; | |
| DAWN_TEST_PARAM_STRUCT(ClampedOOBDynamicBufferOffsetParams, ReadBufferUsage, OOBRead, OOBWrite); | |
| } // anonymous namespace | |
| class ClampedOOBDynamicBufferOffsetTests | |
| : public DawnTestWithParams<ClampedOOBDynamicBufferOffsetParams> {}; | |
| // Test robust buffer access behavior for out of bounds accesses to dynamic buffer bindings. | |
| TEST_P(ClampedOOBDynamicBufferOffsetTests, CheckOOBAccess) { | |
| static constexpr uint32_t kArrayLength = 10u; | |
| // Out-of-bounds access will start halfway into the array and index off the end. | |
| static constexpr uint32_t kOOBOffset = kArrayLength / 2; | |
| wgpu::BufferBindingType sourceBindingType; | |
| switch (GetParam().mReadBufferUsage) { | |
| case wgpu::BufferUsage::Uniform: | |
| sourceBindingType = wgpu::BufferBindingType::Uniform; | |
| break; | |
| case wgpu::BufferUsage::Storage: | |
| sourceBindingType = wgpu::BufferBindingType::ReadOnlyStorage; | |
| break; | |
| default: | |
| DAWN_UNREACHABLE(); | |
| } | |
| wgpu::BindGroupLayout bgl = utils::MakeBindGroupLayout( | |
| device, {{0, wgpu::ShaderStage::Compute, sourceBindingType, true}, | |
| {1, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage, true}}); | |
| wgpu::PipelineLayout layout = utils::MakeBasicPipelineLayout(device, &bgl); | |
| wgpu::ComputePipeline pipeline; | |
| { | |
| std::ostringstream shader; | |
| shader << "const kArrayLength: u32 = " << kArrayLength << "u;\n"; | |
| if (GetParam().mOOBRead) { | |
| shader << "const kReadOffset: u32 = " << kOOBOffset << "u;\n"; | |
| } else { | |
| shader << "const kReadOffset: u32 = 0u;\n"; | |
| } | |
| if (GetParam().mOOBWrite) { | |
| shader << "const kWriteOffset: u32 = " << kOOBOffset << "u;\n"; | |
| } else { | |
| shader << "const kWriteOffset: u32 = 0u;\n"; | |
| } | |
| switch (GetParam().mReadBufferUsage) { | |
| case wgpu::BufferUsage::Uniform: | |
| shader << R"( | |
| struct Src { | |
| values : array<vec4u, kArrayLength> | |
| } | |
| @group(0) @binding(0) var<uniform> src : Src; | |
| )"; | |
| break; | |
| case wgpu::BufferUsage::Storage: | |
| shader << R"( | |
| struct Src { | |
| values : array<vec4u> | |
| } | |
| @group(0) @binding(0) var<storage, read> src : Src; | |
| )"; | |
| break; | |
| default: | |
| DAWN_UNREACHABLE(); | |
| } | |
| shader << R"( | |
| struct Dst { | |
| values : array<vec4u> | |
| } | |
| @group(0) @binding(1) var<storage, read_write> dst : Dst; | |
| )"; | |
| shader << R"( | |
| @compute @workgroup_size(1) fn main() { | |
| for (var i: u32 = 0u; i < kArrayLength; i = i + 1u) { | |
| dst.values[i + kWriteOffset] = src.values[i + kReadOffset]; | |
| } | |
| } | |
| )"; | |
| wgpu::ComputePipelineDescriptor pipelineDesc; | |
| pipelineDesc.layout = layout; | |
| pipelineDesc.compute.module = utils::CreateShaderModule(device, shader.str().c_str()); | |
| pipeline = device.CreateComputePipeline(&pipelineDesc); | |
| } | |
| uint32_t minUniformBufferOffsetAlignment = GetSupportedLimits().minUniformBufferOffsetAlignment; | |
| uint32_t minStorageBufferOffsetAlignment = GetSupportedLimits().minStorageBufferOffsetAlignment; | |
| uint32_t arrayByteLength = kArrayLength * 4 * sizeof(uint32_t); | |
| uint32_t uniformBufferOffset = Align(arrayByteLength, minUniformBufferOffsetAlignment); | |
| uint32_t storageBufferOffset = Align(arrayByteLength, minStorageBufferOffsetAlignment); | |
| // Enough space to bind at a dynamic offset. | |
| uint32_t uniformBufferSize = uniformBufferOffset + arrayByteLength; | |
| uint32_t storageBufferSize = storageBufferOffset + arrayByteLength; | |
| // Buffers are padded so we can check that bytes after the bound range are not changed. | |
| static constexpr uint32_t kEndPadding = 16; | |
| uint64_t srcBufferSize; | |
| uint32_t srcBufferByteOffset; | |
| uint32_t dstBufferByteOffset = storageBufferOffset; | |
| uint64_t dstBufferSize = storageBufferSize + kEndPadding; | |
| switch (GetParam().mReadBufferUsage) { | |
| case wgpu::BufferUsage::Uniform: | |
| srcBufferSize = uniformBufferSize + kEndPadding; | |
| srcBufferByteOffset = uniformBufferOffset; | |
| break; | |
| case wgpu::BufferUsage::Storage: | |
| srcBufferSize = storageBufferSize + kEndPadding; | |
| srcBufferByteOffset = storageBufferOffset; | |
| break; | |
| default: | |
| DAWN_UNREACHABLE(); | |
| } | |
| std::vector<uint32_t> srcData(srcBufferSize / sizeof(uint32_t)); | |
| std::vector<uint32_t> expectedDst(dstBufferSize / sizeof(uint32_t)); | |
| // Fill the src buffer with 0, 1, 2, ... | |
| std::iota(srcData.begin(), srcData.end(), 0); | |
| wgpu::Buffer src = utils::CreateBufferFromData(device, &srcData[0], srcBufferSize, | |
| GetParam().mReadBufferUsage); | |
| // Fill the dst buffer with 0xFF. | |
| DAWN_UNSAFE_TODO(memset(expectedDst.data(), 0xFF, dstBufferSize)); | |
| wgpu::Buffer dst = | |
| utils::CreateBufferFromData(device, &expectedDst[0], dstBufferSize, | |
| wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc); | |
| // Produce expected data assuming the implementation performs clamping. | |
| for (uint32_t i = 0; i < kArrayLength; ++i) { | |
| uint32_t readIndex = GetParam().mOOBRead ? std::min(kOOBOffset + i, kArrayLength - 1) : i; | |
| uint32_t writeIndex = GetParam().mOOBWrite ? std::min(kOOBOffset + i, kArrayLength - 1) : i; | |
| for (uint32_t c = 0; c < 4; ++c) { | |
| uint32_t value = srcData[srcBufferByteOffset / 4 + 4 * readIndex + c]; | |
| expectedDst[dstBufferByteOffset / 4 + 4 * writeIndex + c] = value; | |
| } | |
| } | |
| std::array<uint32_t, 2> dynamicOffsets = {srcBufferByteOffset, dstBufferByteOffset}; | |
| wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, bgl, | |
| { | |
| {0, src, 0, arrayByteLength}, | |
| {1, dst, 0, arrayByteLength}, | |
| }); | |
| wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); | |
| computePassEncoder.SetPipeline(pipeline); | |
| computePassEncoder.SetBindGroup(0, bindGroup, dynamicOffsets.size(), dynamicOffsets.data()); | |
| computePassEncoder.DispatchWorkgroups(1); | |
| computePassEncoder.End(); | |
| wgpu::CommandBuffer commands = commandEncoder.Finish(); | |
| queue.Submit(1, &commands); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expectedDst.data(), dst, 0, dstBufferSize / sizeof(uint32_t)); | |
| } | |
| DAWN_INSTANTIATE_TEST(DynamicBufferOffsetTests, | |
| D3D11Backend(), | |
| D3D12Backend(), | |
| D3D12Backend({}, {"d3d12_use_root_signature_version_1_1"}), | |
| MetalBackend(), | |
| OpenGLBackend(), | |
| OpenGLESBackend(), | |
| VulkanBackend(), | |
| WebGPUBackend()); | |
| // Only instantiate on D3D12 / Metal where we are sure of the robustness implementation. | |
| // Tint injects clamping in the shader. OpenGL(ES) / Vulkan robustness is less constrained. | |
| DAWN_INSTANTIATE_TEST_P(ClampedOOBDynamicBufferOffsetTests, | |
| {D3D12Backend(), D3D12Backend({}, {"d3d12_use_root_signature_version_1_1"}), | |
| MetalBackend()}, | |
| {wgpu::BufferUsage::Uniform, wgpu::BufferUsage::Storage}, | |
| {false, true}, | |
| {false, true}); | |
| class OOBDynamicBufferOffsetTests : public DawnTest {}; | |
| // Tests that a non-visible binding does not end up disabling robustness for a dynamic binding at | |
| // index 0. This was happening on D3D12 because we erroneously were adding non-visible buffers to | |
| // 'ignored_by_robustness' for which the default shader register was 0, thus disabling robustness | |
| // for the dynamic buffer correctly bound to shader register 0. See crbug.com/512988356 | |
| TEST_P(OOBDynamicBufferOffsetTests, NonVisibleBindingDoesNotDisableRobustness) { | |
| // Bind group layout 0: | |
| // binding 0: dynamic uniform, visibility Compute -> robustness should be enabled | |
| // binding 7: read-only-storage, visibility None -> formerly disabled robustness on binding 0 | |
| wgpu::BindGroupLayout bgl0 = utils::MakeBindGroupLayout( | |
| device, { | |
| {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform, true}, | |
| {7, wgpu::ShaderStage::None, wgpu::BufferBindingType::ReadOnlyStorage}, | |
| }); | |
| // Bind group layout 1: index source + output | |
| wgpu::BindGroupLayout bgl1 = utils::MakeBindGroupLayout( | |
| device, { | |
| {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::ReadOnlyStorage}, | |
| {1, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}, | |
| }); | |
| wgpu::PipelineLayout layout = utils::MakePipelineLayout(device, {bgl0, bgl1}); | |
| // WGSL: index a fixed-size uniform array with a runtime value loaded from a | |
| // storage buffer. With robustness enabled, Tint emits `min(idx, 1u)`; with the bug, | |
| // the clamp is erroneously dropped. | |
| wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( | |
| struct U { a : array<u32, 1> }; | |
| @group(0) @binding(0) var<uniform> u : U; | |
| @group(1) @binding(0) var<storage, read> indices : array<u32>; | |
| @group(1) @binding(1) var<storage, read_write> outp: array<u32>; | |
| @compute @workgroup_size(1) | |
| fn main() { | |
| for (var i = 0; i < 8; i++) { | |
| let idx = indices[i]; | |
| let result = u.a[idx]; | |
| outp[i] = result; | |
| } | |
| } | |
| )"); | |
| wgpu::ComputePipelineDescriptor cpDesc; | |
| cpDesc.layout = layout; | |
| cpDesc.compute.module = module; | |
| wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc); | |
| // Uniform with single uint32_t sentinel value | |
| const uint32_t kSentinel = 0xCAFEF00D; | |
| std::vector uniformData(1, kSentinel); | |
| wgpu::Buffer uniformBuf = | |
| utils::CreateBufferFromData(device, uniformData.data(), 1 * sizeof(uint32_t), | |
| wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst); | |
| // Visibility None buffer (must still be bound). | |
| wgpu::BufferDescriptor visNoneDesc; | |
| visNoneDesc.size = sizeof(uint32_t); | |
| visNoneDesc.usage = wgpu::BufferUsage::Storage; | |
| wgpu::Buffer visNoneBuf = device.CreateBuffer(&visNoneDesc); | |
| size_t kBufferElems = 8; | |
| size_t kBufferSize = kBufferElems * sizeof(uint32_t); | |
| // Indices buffer with values 0, 10, 100, 1000, etc. | |
| std::vector<uint32_t> indices(kBufferElems); | |
| for (uint32_t i = 0; i < kBufferElems; ++i) { | |
| indices[i] = i * 10; | |
| } | |
| wgpu::Buffer indicesBuf = | |
| utils::CreateBufferFromData(device, indices.data(), kBufferSize, | |
| wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst); | |
| // Output buffer. | |
| wgpu::BufferDescriptor outDesc; | |
| outDesc.size = kBufferSize; | |
| outDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc; | |
| wgpu::Buffer outBuf = device.CreateBuffer(&outDesc); | |
| wgpu::BindGroup bg0 = utils::MakeBindGroup(device, bgl0, {{0, uniformBuf}, {7, visNoneBuf}}); | |
| wgpu::BindGroup bg1 = utils::MakeBindGroup(device, bgl1, {{0, indicesBuf}, {1, outBuf}}); | |
| wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); | |
| wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); | |
| pass.SetPipeline(pipeline); | |
| uint32_t dynamicOffset = 0; | |
| pass.SetBindGroup(0, bg0, 1, &dynamicOffset); | |
| pass.SetBindGroup(1, bg1); | |
| pass.DispatchWorkgroups(1); | |
| pass.End(); | |
| wgpu::CommandBuffer commands = encoder.Finish(); | |
| device.GetQueue().Submit(1, &commands); | |
| // All values in output should be kSentinel. | |
| // When robustness is erroneously disabled, the uniform buffer is indexed OOB which is UB, so | |
| // the results can be anything, although typically we'll get kSentinel as the first value, and 0 | |
| // for the rest. | |
| std::vector<uint32_t> expected(kBufferSize, kSentinel); | |
| EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), outBuf, 0, kBufferElems); | |
| } | |
| // DAWN_INSTANTIATE_TEST(OOBDynamicBufferOffsetTests, D3D12Backend()); | |
| DAWN_INSTANTIATE_TEST(OOBDynamicBufferOffsetTests, | |
| D3D11Backend(), | |
| D3D12Backend(), | |
| MetalBackend(), | |
| OpenGLBackend(), | |
| OpenGLESBackend(), | |
| VulkanBackend(), | |
| WebGPUBackend()); | |
| } // anonymous namespace | |
| } // namespace dawn |