// Copyright 2024 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 <cstdint>
#include <limits>
#include <string>
#include <vector>

#include "dawn/common/NonMovable.h"
#include "dawn/tests/unittests/validation/ValidationTest.h"
#include "dawn/utils/ComboRenderPipelineDescriptor.h"
#include "dawn/utils/WGPUHelpers.h"

namespace dawn {
namespace {

enum class FeatureMode {
    Enabled,
    DisabledViaNotAllowUnsafeAPIs,
    DisabledViaBlocklistedFeatures,
};

// Test that the feature only works when enabled
struct ImmediateDataDisableTest : ValidationTestWithParam<FeatureMode> {
    std::vector<const char*> GetWGSLBlocklistedFeatures() override {
        switch (GetParam()) {
            case FeatureMode::Enabled:
                return {};
            case FeatureMode::DisabledViaNotAllowUnsafeAPIs:
                return {};
            case FeatureMode::DisabledViaBlocklistedFeatures:
                return {"immediate_address_space"};
        }
        DAWN_UNREACHABLE();
        return {};
    }

    bool AllowUnsafeAPIs() override {
        switch (GetParam()) {
            case FeatureMode::Enabled:
                // Currently the only way to enable ImmediateAddressSpace is via AllowUnsafeAPIs.
                // See GetLanguageFeatureStatus.
                return true;
            case FeatureMode::DisabledViaNotAllowUnsafeAPIs:
                return false;
            case FeatureMode::DisabledViaBlocklistedFeatures:
                // Enabling AllowUnsafeAPIs while disabling via blocklist should still fail.
                return true;
        }
        DAWN_UNREACHABLE();
        return false;
    }
};

// Check that creating a PipelineLayout with non-zero immediateSize is disallowed
// without the feature enabled.
TEST_P(ImmediateDataDisableTest, ImmediateSizeNotAllowed) {
    wgpu::PipelineLayoutDescriptor desc{};
    desc.immediateSize = 1;

    if (GetParam() == FeatureMode::Enabled) {
        device.CreatePipelineLayout(&desc);
    } else {
        ASSERT_DEVICE_ERROR(device.CreatePipelineLayout(&desc));
    }
}

// Check that SetImmediates doesn't work (even with size=0) without the feature enabled.
TEST_P(ImmediateDataDisableTest, SetImmediates) {
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
        pass.SetImmediates(0, nullptr, 0);
        pass.End();
        if (GetParam() == FeatureMode::Enabled) {
            encoder.Finish();
        } else {
            ASSERT_DEVICE_ERROR(encoder.Finish());
        }
    }
    {
        const uint32_t data = 0;

        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
        pass.SetImmediates(0, &data, 4);
        pass.End();
        if (GetParam() == FeatureMode::Enabled) {
            encoder.Finish();
        } else {
            ASSERT_DEVICE_ERROR(encoder.Finish());
        }
    }
}

// Check that limits.maxImmediateSize is 0 when the feature is disabled, and kMaxImmediateDataBytes
// otherwise.
TEST_P(ImmediateDataDisableTest, MaxImmediateSizeIsZero) {
    if (GetParam() == FeatureMode::Enabled) {
        ASSERT_EQ(GetSupportedLimits().maxImmediateSize, kMaxImmediateDataBytes);
    } else {
        ASSERT_EQ(GetSupportedLimits().maxImmediateSize, 0u);
    }
}

INSTANTIATE_TEST_SUITE_P(,
                         ImmediateDataDisableTest,
                         ::testing::ValuesIn({FeatureMode::Enabled,
                                              FeatureMode::DisabledViaNotAllowUnsafeAPIs,
                                              FeatureMode::DisabledViaBlocklistedFeatures}));

class ImmediateDataTest : public ValidationTest {
  protected:
    wgpu::BindGroupLayout CreateBindGroupLayout() {
        wgpu::BindGroupLayoutEntry entries[1];
        entries[0].binding = 0;
        entries[0].visibility = wgpu::ShaderStage::Compute;
        entries[0].buffer.type = wgpu::BufferBindingType::Storage;

        wgpu::BindGroupLayoutDescriptor bindGroupLayoutDesc;
        bindGroupLayoutDesc.entryCount = 1;
        bindGroupLayoutDesc.entries = entries;

        return device.CreateBindGroupLayout(&bindGroupLayoutDesc);
    }

    wgpu::PipelineLayout CreatePipelineLayout(uint32_t requiredImmediateSize) {
        wgpu::BindGroupLayout bindGroupLayout = CreateBindGroupLayout();

        wgpu::PipelineLayoutDescriptor pipelineLayoutDesc;
        pipelineLayoutDesc.bindGroupLayoutCount = 1;
        pipelineLayoutDesc.bindGroupLayouts = &bindGroupLayout;
        pipelineLayoutDesc.immediateSize = requiredImmediateSize;
        return device.CreatePipelineLayout(&pipelineLayoutDesc);
    }

    wgpu::ShaderModule mShaderModule;
};

// Check that non-zero immediateSize is possible with feature enabled and size must
// below max size limits.
TEST_F(ImmediateDataTest, ValidateImmediateSize) {
    wgpu::PipelineLayoutDescriptor desc{};

    // Success case with valid immediateSize.
    {
        desc.immediateSize = kMaxImmediateDataBytes;
        device.CreatePipelineLayout(&desc);
    }

    // Failed case with invalid immediateSize that exceed limits.
    {
        desc.immediateSize = kMaxImmediateDataBytes + 1;
        ASSERT_DEVICE_ERROR(device.CreatePipelineLayout(&desc));
    }
}

// Check that SetImmediates offset and length must be aligned to 4 bytes.
TEST_F(ImmediateDataTest, ValidateSetImmediatesAlignment) {
    // Success cases
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t data = 0;
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(0, &data, 4);
        computePass.End();
        encoder.Finish();
    }

    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(4, nullptr, 0);
        computePass.End();
        encoder.Finish();
    }

    // Failed case with non-aligned offset bytes
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(2, nullptr, 0);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }

    // Failed cases with non-aligned size
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint8_t data = 0;
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(0, &data, 2);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }
}

// Check that SetImmediates offset + length must be in bound.
TEST_F(ImmediateDataTest, ValidateSetImmediatesOOB) {
    // Success cases
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        std::vector<uint32_t> data(kMaxImmediateDataBytes / 4, 0);
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(0, data.data(), kMaxImmediateDataBytes);
        computePass.End();
        encoder.Finish();
    }

    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(kMaxImmediateDataBytes, nullptr, 0);
        computePass.End();
        encoder.Finish();
    }

    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t data = 0;
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(kMaxImmediateDataBytes - 4, &data, 4);
        computePass.End();
        encoder.Finish();
    }

    // Failed case with offset oob
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t offset = kMaxImmediateDataBytes + 4;
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(offset, nullptr, 0);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }

    // Failed cases with size oob
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t size = kMaxImmediateDataBytes + 4;
        std::vector<uint32_t> data(size / 4, 0);
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(0, data.data(), size);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }

    // Failed cases with offset + size oob
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t offset = kMaxImmediateDataBytes;
        uint32_t data[] = {0};
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(offset, data, 4);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }

    // Failed case with super large offset + size oob but looping back to zero
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        uint32_t offset = std::numeric_limits<uint32_t>::max() - 3;
        uint32_t data[] = {0};
        wgpu::ComputePassEncoder computePass = encoder.BeginComputePass();
        computePass.SetImmediates(offset, data, 4);
        computePass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }
}

// Check that pipelineLayout immediate data bytes compatible with shaders.
TEST_F(ImmediateDataTest, ValidatePipelineLayoutImmediateDataBytesAndShaders) {
    constexpr uint32_t kShaderImmediateDataBytes = 12u;
    wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, R"(
        var<immediate> fragmentConstants: vec3f;
        var<immediate> computeConstants: vec3u;
        @vertex fn vsMain(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f {
            const pos = array(
                vec2( 1.0, -1.0),
                vec2(-1.0, -1.0),
                vec2( 0.0,  1.0),
            );
            return vec4(pos[VertexIndex], 0.0, 1.0);
        }

        // to reuse the same pipeline layout
        @fragment fn fsMain() -> @location(0) vec4f {
            return vec4f(fragmentConstants, 1.0);
        }


        @group(0) @binding(0) var<storage, read_write> output : vec3u;

        @compute @workgroup_size(1, 1, 1)
        fn csMain() {
            output = computeConstants;
        })");

    // Success cases
    {
        utils::ComboRenderPipelineDescriptor pipelineDescriptor;
        pipelineDescriptor.vertex.module = shaderModule;
        pipelineDescriptor.cFragment.module = shaderModule;
        pipelineDescriptor.layout = CreatePipelineLayout(kShaderImmediateDataBytes);
        device.CreateRenderPipeline(&pipelineDescriptor);
    }

    {
        wgpu::ComputePipelineDescriptor csDesc;
        csDesc.compute.module = shaderModule;
        csDesc.layout = CreatePipelineLayout(kShaderImmediateDataBytes);

        device.CreateComputePipeline(&csDesc);
    }

    // Default layout
    {
        utils::ComboRenderPipelineDescriptor pipelineDescriptor;
        pipelineDescriptor.vertex.module = shaderModule;
        pipelineDescriptor.cFragment.module = shaderModule;
        device.CreateRenderPipeline(&pipelineDescriptor);
    }

    {
        wgpu::ComputePipelineDescriptor csDesc;
        csDesc.compute.module = shaderModule;

        device.CreateComputePipeline(&csDesc);
    }

    // Failed case with fragment shader requires more immediate data.
    {
        utils::ComboRenderPipelineDescriptor pipelineDescriptor;
        pipelineDescriptor.vertex.module = shaderModule;
        pipelineDescriptor.cFragment.module = shaderModule;
        pipelineDescriptor.layout = CreatePipelineLayout(kShaderImmediateDataBytes - 4);
        ASSERT_DEVICE_ERROR(device.CreateRenderPipeline(&pipelineDescriptor));
    }

    // Failed cases with compute shader requires more immediate data.
    {
        wgpu::ComputePipelineDescriptor csDesc;
        csDesc.compute.module = shaderModule;
        csDesc.layout = CreatePipelineLayout(kShaderImmediateDataBytes - 4);

        ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
    }
}

// Check that default pipelineLayout has too many immediate data bytes .
TEST_F(ImmediateDataTest, ValidateDefaultPipelineLayout) {
    wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, R"(
            var<immediate> fragmentConstants: vec4f;

            var<immediate> computeConstants: vec4u;

            @vertex fn vsMain(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f {
                const pos = array(
                    vec2( 1.0, -1.0),
                    vec2(-1.0, -1.0),
                    vec2( 0.0,  1.0),
                );
                return vec4(pos[VertexIndex], 0.0, 1.0);
            }

            // to reuse the same pipeline layout
            @fragment fn fsMain() -> @location(0) vec4f {
                return vec4f(fragmentConstants.x, fragmentConstants.yzw);
            }

            @group(0) @binding(0) var<storage, read_write> output : vec4u;

            @compute @workgroup_size(1, 1, 1)
            fn csMain() {
                output = vec4u(computeConstants.x, computeConstants.yzw);
            })");

    // Failed case with too much immediate data in shader
    ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
            struct FragmentConstants {
                c0: vec4f,
                c1: vec4f,
                c2: vec4f,
                c3: vec4f,
                constantsOOB: f32,
            };

            struct ComputeConstants {
                c0: vec4u,
                c1: vec4u,
                c2: vec4u,
                c3: vec4u,
                constantsOOB: u32,
            };
            var<immediate> fragmentConstants: FragmentConstants;
            var<immediate> computeConstants: ComputeConstants;
            @vertex fn vsMain(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f {
                const pos = array(
                    vec2( 1.0, -1.0),
                    vec2(-1.0, -1.0),
                    vec2( 0.0,  1.0),
                );
                return vec4(pos[VertexIndex], 0.0, 1.0);
            }

            // to reuse the same pipeline layout
            @fragment fn fsMain() -> @location(0) vec4f {
                return vec4f(fragmentConstants.c0.x + fragmentConstants.constantsOOB,
                             fragmentConstants.c0.yzw);
            }

            @group(0) @binding(0) var<storage, read_write> output : vec4u;

            @compute @workgroup_size(1, 1, 1)
            fn csMain() {
                output = vec4u(computeConstants.c0.x + computeConstants.constantsOOB,
                               computeConstants.c0.yzw);
            })"));

    // Success cases
    {
        utils::ComboRenderPipelineDescriptor pipelineDescriptor;
        pipelineDescriptor.vertex.module = shaderModule;
        pipelineDescriptor.cFragment.module = shaderModule;
        device.CreateRenderPipeline(&pipelineDescriptor);
    }

    {
        wgpu::ComputePipelineDescriptor csDesc;
        csDesc.compute.module = shaderModule;

        device.CreateComputePipeline(&csDesc);
    }
}

// Check that executing multiple bundles with different immediate data requirements works.
TEST_F(ImmediateDataTest, ExecuteBundlesWithDifferentImmediateData) {
    // Pipeline 4: requires 4 bytes
    wgpu::ShaderModule module4 = utils::CreateShaderModule(device, R"(
        var<immediate> constants: u32;
        @vertex fn vs() -> @builtin(position) vec4f {
            _ = constants;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
        @fragment fn fs() -> @location(0) vec4f {
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
    )");
    utils::ComboRenderPipelineDescriptor desc4;
    desc4.vertex.module = module4;
    desc4.cFragment.module = module4;
    wgpu::RenderPipeline pipeline4 = device.CreateRenderPipeline(&desc4);

    // Pipeline 8: requires 8 bytes
    wgpu::ShaderModule module8 = utils::CreateShaderModule(device, R"(
        struct Constants {
            a: u32,
            b: u32,
        }
        var<immediate> constants: Constants;
        @vertex fn vs() -> @builtin(position) vec4f {
            _ = constants.b;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
        @fragment fn fs() -> @location(0) vec4f {
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
    )");
    utils::ComboRenderPipelineDescriptor desc8;
    desc8.vertex.module = module8;
    desc8.cFragment.module = module8;
    wgpu::RenderPipeline pipeline8 = device.CreateRenderPipeline(&desc8);

    utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);

    // Bundle 4
    wgpu::RenderBundle bundle4;
    {
        wgpu::RenderBundleEncoderDescriptor bundleDesc;
        bundleDesc.colorFormatCount = 1;
        bundleDesc.colorFormats = &renderPass.colorFormat;
        wgpu::RenderBundleEncoder encoder = device.CreateRenderBundleEncoder(&bundleDesc);
        encoder.SetPipeline(pipeline4);
        uint32_t data = 0;
        encoder.SetImmediates(0, &data, 4);
        encoder.Draw(3);
        bundle4 = encoder.Finish();
    }

    // Bundle 8
    wgpu::RenderBundle bundle8;
    {
        wgpu::RenderBundleEncoderDescriptor bundleDesc;
        bundleDesc.colorFormatCount = 1;
        bundleDesc.colorFormats = &renderPass.colorFormat;
        wgpu::RenderBundleEncoder encoder = device.CreateRenderBundleEncoder(&bundleDesc);
        encoder.SetPipeline(pipeline8);
        uint32_t data[] = {0, 0};
        encoder.SetImmediates(0, data, 8);
        encoder.Draw(3);
        bundle8 = encoder.Finish();
    }

    // Execute both
    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
    wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
    wgpu::RenderBundle bundles[] = {bundle4, bundle8};
    pass.ExecuteBundles(2, bundles);
    pass.End();
    encoder.Finish();
}

// Check that ExecuteBundles resets the immediate data state in the RenderPass.
TEST_F(ImmediateDataTest, ExecuteBundlesResetsImmediateDataState) {
    // Pipeline 4: requires 4 bytes
    wgpu::ShaderModule module4 = utils::CreateShaderModule(device, R"(
        var<immediate> constants: u32;
        @vertex fn vs() -> @builtin(position) vec4f {
            _ = constants;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
        @fragment fn fs() -> @location(0) vec4f {
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
    )");
    utils::ComboRenderPipelineDescriptor desc4;
    desc4.vertex.module = module4;
    desc4.cFragment.module = module4;
    wgpu::RenderPipeline pipeline4 = device.CreateRenderPipeline(&desc4);

    // Bundle (placeholder, just to execute)
    wgpu::RenderBundle bundle;
    {
        wgpu::RenderBundleEncoderDescriptor bundleDesc;
        bundleDesc.colorFormatCount = 1;
        wgpu::TextureFormat format = wgpu::TextureFormat::RGBA8Unorm;
        bundleDesc.colorFormats = &format;
        wgpu::RenderBundleEncoder encoder = device.CreateRenderBundleEncoder(&bundleDesc);
        bundle = encoder.Finish();
    }

    // Case 1: Immediate -> ExecuteBundles -> SetPipeline -> Draw (Fail: No immediate data)
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);
        wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);

        pass.SetPipeline(pipeline4);
        uint32_t data = 0;
        pass.SetImmediates(0, &data, 4);

        pass.ExecuteBundles(1, &bundle);

        pass.SetPipeline(pipeline4);  // Restore pipeline
        pass.Draw(3);                 // Should fail (Immediate data lost)
        pass.End();
        ASSERT_DEVICE_ERROR(encoder.Finish());
    }

    // Case 2: ExecuteBundles -> SetImmediates -> SetPipeline -> Draw (Success)
    {
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
        utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);
        wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
        pass.ExecuteBundles(1, &bundle);
        uint32_t data = 0;
        pass.SetImmediates(0, &data, 4);
        pass.SetPipeline(pipeline4);
        pass.Draw(3);
        pass.End();
        encoder.Finish();
    }
}

enum class EncoderType {
    Compute,
    RenderPass,
    RenderBundle,
};

struct ImmediateDataRange {
    uint32_t offset;
    uint32_t size;
};

class ImmediateDataRequiredTest : public ImmediateDataTest,
                                  public testing::WithParamInterface<EncoderType> {
  protected:
    void TestImmediateDataValidation(std::string entryPoint,
                                     std::vector<ImmediateDataRange> ranges,
                                     bool success,
                                     EncoderType encoderType) {
        // Structs with padding:
        // PadMiddle:
        //   a: u32 (0..4)
        //   padding (4..16) -> 12 bytes
        //   b: vec4<f32> (16..32)
        //   Size: 32
        //
        // PadTail:
        //   a: vec3<f32> (0..12)
        //   padding (12..16) -> 4 bytes
        //   Size: 16
        //
        // Layout:
        //   padMiddle: offset 0, size 32
        //   padTail: offset 32, size 16
        const char* kShader = R"(
        struct PadMiddle {
            a : u32,
            b : vec4<f32>,
        }
        struct PadTail {
            a : vec3<f32>,
        }
        var<immediate> padMiddle : PadMiddle;
        var<immediate> padTail : PadTail;

        @compute @workgroup_size(1)
        fn mainMiddle() {
            _ = padMiddle.b;
        }

        @compute @workgroup_size(1)
        fn mainTail() {
            _ = padTail.a;
        }

        @vertex fn vsMiddle() -> @builtin(position) vec4f {
            _ = padMiddle.b;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
        @fragment fn fsMiddle() -> @location(0) vec4f {
            _ = padMiddle.b;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }

        @vertex fn vsTail() -> @builtin(position) vec4f {
            _ = padTail.a;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
        @fragment fn fsTail() -> @location(0) vec4f {
            _ = padTail.a;
            return vec4f(0.0, 0.0, 0.0, 1.0);
        }
    )";

        wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, kShader);
        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();

        auto SetImmediates = [&](auto& pass) {
            for (const auto& range : ranges) {
                if (range.size > 0) {
                    std::vector<uint32_t> data((range.size + 3) / 4, 0);
                    pass.SetImmediates(range.offset, data.data(), range.size);
                }
            }
        };

        auto CreateRenderPipeline = [&]() {
            utils::ComboRenderPipelineDescriptor descriptor;
            descriptor.vertex.module = shaderModule;
            descriptor.cFragment.module = shaderModule;

            if (entryPoint == "mainMiddle") {
                descriptor.vertex.entryPoint = "vsMiddle";
                descriptor.cFragment.entryPoint = "fsMiddle";
            } else {
                descriptor.vertex.entryPoint = "vsTail";
                descriptor.cFragment.entryPoint = "fsTail";
            }

            return device.CreateRenderPipeline(&descriptor);
        };

        switch (encoderType) {
            case EncoderType::Compute: {
                wgpu::ComputePipelineDescriptor descriptor;
                descriptor.compute.module = shaderModule;
                descriptor.compute.entryPoint = entryPoint.c_str();
                wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&descriptor);

                wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
                pass.SetPipeline(pipeline);
                SetImmediates(pass);
                pass.DispatchWorkgroups(1);
                pass.End();
                break;
            }
            case EncoderType::RenderPass: {
                wgpu::RenderPipeline pipeline = CreateRenderPipeline();
                utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1);
                wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo);
                pass.SetPipeline(pipeline);
                SetImmediates(pass);
                pass.Draw(3);
                pass.End();
                break;
            }
            case EncoderType::RenderBundle: {
                wgpu::RenderPipeline pipeline = CreateRenderPipeline();
                wgpu::RenderBundleEncoderDescriptor bundleDesc;
                bundleDesc.colorFormatCount = 1;
                wgpu::TextureFormat format = wgpu::TextureFormat::RGBA8Unorm;
                bundleDesc.colorFormats = &format;

                wgpu::RenderBundleEncoder bundleEncoder =
                    device.CreateRenderBundleEncoder(&bundleDesc);
                bundleEncoder.SetPipeline(pipeline);
                SetImmediates(bundleEncoder);
                bundleEncoder.Draw(3);

                if (success) {
                    bundleEncoder.Finish();
                } else {
                    ASSERT_DEVICE_ERROR(bundleEncoder.Finish());
                }
                return;
            }
        }

        if (success) {
            encoder.Finish();
        } else {
            ASSERT_DEVICE_ERROR(encoder.Finish());
        }
    }

    void RunTests(std::string entryPoint, std::vector<ImmediateDataRange> ranges, bool success) {
        TestImmediateDataValidation(entryPoint, ranges, success, GetParam());
    }
};

TEST_P(ImmediateDataRequiredTest, PadMiddleMissesA) {
    RunTests("mainMiddle", {{16, 16}}, false);
}

TEST_P(ImmediateDataRequiredTest, PadMiddleCoversAll) {
    RunTests("mainMiddle", {{0, 32}}, true);
}

TEST_P(ImmediateDataRequiredTest, PadMiddleMissesB) {
    RunTests("mainMiddle", {{0, 16}}, false);
}

TEST_P(ImmediateDataRequiredTest, PadMiddlePartialB) {
    RunTests("mainMiddle", {{16, 12}}, false);
}

TEST_P(ImmediateDataRequiredTest, PadMiddleSplitCoverage) {
    RunTests("mainMiddle", {{0, 4}, {16, 16}}, true);
}

TEST_P(ImmediateDataRequiredTest, PadTailCoversA) {
    RunTests("mainTail", {{0, 12}}, true);
}

TEST_P(ImmediateDataRequiredTest, PadTailCoversAll) {
    RunTests("mainTail", {{0, 16}}, true);
}

TEST_P(ImmediateDataRequiredTest, PadTailPartialA) {
    RunTests("mainTail", {{0, 8}}, false);
}

INSTANTIATE_TEST_SUITE_P(
    ,
    ImmediateDataRequiredTest,
    ::testing::Values(EncoderType::Compute, EncoderType::RenderPass, EncoderType::RenderBundle),
    [](const testing::TestParamInfo<ImmediateDataRequiredTest::ParamType>& info) {
        switch (info.param) {
            case EncoderType::Compute:
                return "Compute";
            case EncoderType::RenderPass:
                return "RenderPass";
            case EncoderType::RenderBundle:
                return "RenderBundle";
        }
        return "Unknown";
    });

}  // anonymous namespace
}  // namespace dawn