// 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 =
        static_cast<uint32_t>(static_cast<size_t>(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