// Copyright 2026 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 <initializer_list>
#include <string>
#include <vector>

#include "dawn/common/Math.h"
#include "dawn/tests/DawnTest.h"
#include "dawn/utils/WGPUHelpers.h"

namespace dawn {
namespace {

class BufferViewTest : public DawnTest {
  protected:
    void SetUp() override { DawnTest::SetUp(); }

    void BaseLengthTest(const std::string& wgsl,
                        uint32_t len,
                        const std::initializer_list<uint32_t>& uniforms,
                        const std::vector<uint32_t>& expected) {
        wgpu::ShaderModule module = utils::CreateShaderModule(device, wgsl);

        std::vector<wgpu::ConstantEntry> constants{{nullptr, "len", static_cast<double>(len)}};
        wgpu::ComputePipelineDescriptor desc;
        desc.compute.module = module;
        desc.compute.constants = constants.data();
        desc.compute.constantCount = constants.size();
        wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&desc);

        wgpu::BufferDescriptor srcDesc;
        srcDesc.size = len;
        srcDesc.usage =
            wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
        wgpu::Buffer src = device.CreateBuffer(&srcDesc);

        wgpu::BufferDescriptor outDesc;
        outDesc.size = sizeof(uint32_t) * expected.size();
        outDesc.usage =
            wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
        wgpu::Buffer out = device.CreateBuffer(&outDesc);

        wgpu::Buffer uni = utils::CreateBufferFromData<uint32_t>(
            device, wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopySrc, uniforms);

        wgpu::BindGroup bg =
            utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
                                 {
                                     {0, src, 0, srcDesc.size},
                                     {1, out, 0, outDesc.size},
                                     {2, uni, 0, uniforms.size() * sizeof(uint32_t)},
                                 });

        wgpu::CommandBuffer commands;
        {
            wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
            wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
            pass.SetPipeline(pipeline);
            pass.SetBindGroup(0, bg);
            pass.DispatchWorkgroups(1, 1, 1);
            pass.End();
            commands = encoder.Finish();
        }
        queue.Submit(1, &commands);

        EXPECT_BUFFER_U32_RANGE_EQ(&expected[0], out, 0, expected.size())
            << "Buffer length = " << len << "\n";
    }

    void TestBufferViewArrayLength(uint32_t len) {
        const std::string src = R"(
          override len : u32;
          struct Offsets {
            a : u32,
            b : u32,
            c : u32,
          }

          @group(0) @binding(0) var<storage> in : buffer;
          @group(0) @binding(1) var<storage, read_write> out : array<u32>;
          @group(0) @binding(2) var<uniform> offsets : Offsets;

          struct S1 {
            a : vec4f,
            b : array<u32>,
          }
          struct S2 {
            a : vec4f,
            b : array<vec2u>,
          }
          struct S3 {
            a : vec4f,
            b : array<vec4u>,
          }

          @compute @workgroup_size(1)
          fn main() {
            out[0] = bufferLength(&in);

            out[1] = arrayLength(bufferView<array<u32>>(&in, offsets.a));
            out[2] = arrayLength(bufferView<array<u32>>(&in, offsets.b));
            out[3] = arrayLength(bufferView<array<u32>>(&in, offsets.c));

            out[4] = arrayLength(bufferView<array<vec2u>>(&in, offsets.a));
            out[5] = arrayLength(bufferView<array<vec2u>>(&in, offsets.b));
            out[6] = arrayLength(bufferView<array<vec2u>>(&in, offsets.c));

            out[7] = arrayLength(bufferView<array<vec3u>>(&in, offsets.a));
            out[8] = arrayLength(bufferView<array<vec3u>>(&in, offsets.b));
            out[9] = arrayLength(bufferView<array<vec3u>>(&in, offsets.c));

            out[10] = arrayLength(bufferView<array<vec4u>>(&in, offsets.a));
            out[11] = arrayLength(bufferView<array<vec4u>>(&in, offsets.b));
            out[12] = arrayLength(bufferView<array<vec4u>>(&in, offsets.c));

            out[13] = arrayLength(&bufferView<S1>(&in, offsets.a).b);
            out[14] = arrayLength(&bufferView<S1>(&in, offsets.b).b);
            out[15] = arrayLength(&bufferView<S1>(&in, offsets.c).b);

            out[16] = arrayLength(&bufferView<S2>(&in, offsets.a).b);
            out[17] = arrayLength(&bufferView<S2>(&in, offsets.b).b);
            out[18] = arrayLength(&bufferView<S2>(&in, offsets.c).b);

            out[19] = arrayLength(&bufferView<S3>(&in, offsets.a).b);
            out[20] = arrayLength(&bufferView<S3>(&in, offsets.b).b);
            out[21] = arrayLength(&bufferView<S3>(&in, offsets.c).b);
          }
        )";

        std::initializer_list<uint32_t> uniforms = {0u, 16u, 32u};

        std::vector<uint32_t> expected = {
            len,                   //
            len / 4,               //
            (len - 16) / 4,        //
            (len - 32) / 4,        //
            len / 8,               //
            (len - 16) / 8,        //
            (len - 32) / 8,        //
            len / 16,              //
            (len - 16) / 16,       //
            (len - 32) / 16,       //
            len / 16,              //
            (len - 16) / 16,       //
            (len - 32) / 16,       //
            (len - 16) / 4,        //
            (len - 16 - 16) / 4,   //
            (len - 16 - 32) / 4,   //
            (len - 16) / 8,        //
            (len - 16 - 16) / 8,   //
            (len - 16 - 32) / 8,   //
            (len - 16) / 16,       //
            (len - 16 - 16) / 16,  //
            (len - 16 - 32) / 16,  //
        };

        BaseLengthTest(src, len, uniforms, expected);
    }

    void TestBufferArrayViewArrayLength(uint32_t len) {
        const std::string src = R"(
          override len : u32;
          @group(0) @binding(0) var<storage> in : buffer;
          @group(0) @binding(1) var<storage, read_write> out : array<u32>;
          @group(0) @binding(2) var<uniform> uniforms : array<vec4u, 2>;

          struct S1 {
            a : u32,
            b : array<u32>,
          }
          struct S2 {
            a : u32,
            b : array<vec2u>,
          }
          struct S3 {
            a : u32,
            b : array<vec4u>,
          }

          @compute @workgroup_size(1)
          fn main() {
            var i = 0u;
            out[i] = bufferLength(&in); i++;

            out[i] = arrayLength(bufferArrayView<array<u32>>(&in, uniforms[0][0], uniforms[1][0])); i++;
            out[i] = arrayLength(bufferArrayView<array<u32>>(&in, uniforms[0][1], uniforms[1][1])); i++;
            out[i] = arrayLength(bufferArrayView<array<u32>>(&in, uniforms[0][2], uniforms[1][2])); i++;
            out[i] = arrayLength(bufferArrayView<array<u32>>(&in, uniforms[0][3], uniforms[1][3])); i++;

            out[i] = arrayLength(bufferArrayView<array<vec2u>>(&in, uniforms[0][0], uniforms[1][0])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec2u>>(&in, uniforms[0][1], uniforms[1][1])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec2u>>(&in, uniforms[0][2], uniforms[1][2])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec2u>>(&in, uniforms[0][2], uniforms[1][3])); i++;

            out[i] = arrayLength(bufferArrayView<array<vec3u>>(&in, uniforms[0][0], uniforms[1][0])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec3u>>(&in, uniforms[0][1], uniforms[1][1])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec3u>>(&in, uniforms[0][2], uniforms[1][2])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec3u>>(&in, uniforms[0][1], uniforms[1][3])); i++;

            out[i] = arrayLength(bufferArrayView<array<vec4u>>(&in, uniforms[0][0], uniforms[1][0])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec4u>>(&in, uniforms[0][1], uniforms[1][1])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec4u>>(&in, uniforms[0][2], uniforms[1][2])); i++;
            out[i] = arrayLength(bufferArrayView<array<vec4u>>(&in, uniforms[0][0], uniforms[1][3])); i++;

            out[i] = arrayLength(&bufferArrayView<S1>(&in, uniforms[0][0], uniforms[1][0]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S1>(&in, uniforms[0][1], uniforms[1][1]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S1>(&in, uniforms[0][2], uniforms[1][2]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S1>(&in, uniforms[0][1], uniforms[1][3]).b); i++;

            out[i] = arrayLength(&bufferArrayView<S2>(&in, uniforms[0][0], uniforms[1][0]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S2>(&in, uniforms[0][1], uniforms[1][1]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S2>(&in, uniforms[0][2], uniforms[1][2]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S2>(&in, uniforms[0][1], uniforms[1][3]).b); i++;

            out[i] = arrayLength(&bufferArrayView<S3>(&in, uniforms[0][0], uniforms[1][0]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S3>(&in, uniforms[0][1], uniforms[1][1]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S3>(&in, uniforms[0][2], uniforms[1][2]).b); i++;
            out[i] = arrayLength(&bufferArrayView<S3>(&in, uniforms[0][0], uniforms[1][3]).b); i++;
          }
        )";

        std::initializer_list<uint32_t> uniforms = {0u, 16u, 32u, 48u, 20u, 36u, 52u, len - 20u};
        std::vector<uint32_t> uniformData{uniforms};

        auto InBounds = [&](uint32_t offset_idx, uint32_t size_idx, uint32_t size,
                            uint32_t str_offset = 0) {
            auto req_size = uniformData[size_idx];
            auto used_size = std::max(((req_size - str_offset) / size) * size, size);
            if (len < str_offset + uniformData[offset_idx] + used_size) {
                return 1u;
            }
            return used_size / size;
        };

        std::vector<uint32_t> expected = {
            len,                         //
            InBounds(0, 4 + 0, 4),       //
            InBounds(1, 4 + 1, 4),       //
            InBounds(2, 4 + 2, 4),       //
            InBounds(3, 4 + 3, 4),       //
            InBounds(0, 4 + 0, 8),       //
            InBounds(1, 4 + 1, 8),       //
            InBounds(2, 4 + 2, 8),       //
            InBounds(2, 4 + 3, 8),       //
            InBounds(0, 4 + 0, 16),      //
            InBounds(1, 4 + 1, 16),      //
            InBounds(2, 4 + 2, 16),      //
            InBounds(1, 4 + 3, 16),      //
            InBounds(0, 4 + 0, 16),      //
            InBounds(1, 4 + 1, 16),      //
            InBounds(2, 4 + 2, 16),      //
            InBounds(0, 4 + 3, 16),      //
            InBounds(0, 4 + 0, 4, 4),    //
            InBounds(1, 4 + 1, 4, 4),    //
            InBounds(2, 4 + 2, 4, 4),    //
            InBounds(1, 4 + 3, 4, 4),    //
            InBounds(0, 4 + 0, 8, 8),    //
            InBounds(1, 4 + 1, 8, 8),    //
            InBounds(2, 4 + 2, 8, 8),    //
            InBounds(1, 4 + 3, 8, 8),    //
            InBounds(0, 4 + 0, 16, 16),  //
            InBounds(1, 4 + 1, 16, 16),  //
            InBounds(2, 4 + 2, 16, 16),  //
            InBounds(0, 4 + 3, 16, 16),  //
        };

        BaseLengthTest(src, len, uniforms, expected);
    }

    void TestArrayLengthFunctions(uint32_t len) {
        const std::string wgsl = R"(
          override len : u32;
          @group(0) @binding(0) var<storage> in : buffer<)" +
                                 std::to_string(len) + R"(>;
          @group(0) @binding(1) var<storage, read_write> out : array<u32>;
          @group(0) @binding(2) var<uniform> uniforms : buffer<20>;

          var<workgroup> wg1 : buffer<64>;
          var<workgroup> wg2 : buffer<128>;
          var<workgroup> wg3 : buffer<len>;

          struct S {
            a : vec2u,
            b : array<u32>,
          }

          fn uniformLengthUnsized(p : ptr<uniform, buffer>, offset : u32) -> u32 {
            return arrayLength(bufferView<array<u32>>(p, offset));
          }

          fn uniformLengthSized(p : ptr<uniform, buffer<16>>, offset : u32) -> u32 {
            return uniformLengthUnsized(p, offset);
          }

          fn workgroupLengthUnsized(p : ptr<workgroup, buffer>, offset : u32, size : u32) -> u32 {
            return arrayLength(bufferArrayView<array<u32>>(p, offset, size));
          }

          fn workgroupLengthSmall(p : ptr<workgroup, buffer<64>>, offset : u32) -> u32 {
            return workgroupLengthUnsized(p, offset, 32);
          }

          fn workgroupLengthOverride(p : ptr<workgroup, buffer<len>>, offset : u32) -> u32 {
            return workgroupLengthUnsized(p, offset, 32);
          }

          fn storageLengthUnsized(p : ptr<storage, buffer>, offset : u32, size : u32) -> u32 {
            return arrayLength(&bufferArrayView<S>(p, offset, size).b);
          }

          fn storageLengthLarge(p : ptr<storage, buffer>, offset : u32) -> u32 {
            return arrayLength(&bufferView<S>(p, offset).b);
          }

          fn storageLengthSmall(p : ptr<storage, buffer<48>>, offset : u32) -> u32 {
            return storageLengthLarge(p, offset);
          }

          @compute @workgroup_size(1)
          fn main() {
            let offsets = bufferView<array<u32>>(&uniforms, 0u);
            var i = 0u;
            out[i] = uniformLengthUnsized(&uniforms, offsets[0]); i++;
            out[i] = uniformLengthUnsized(&uniforms, offsets[1]); i++;
            out[i] = uniformLengthUnsized(&uniforms, offsets[2]); i++;
            out[i] = uniformLengthSized(&uniforms, offsets[0]); i++;
            out[i] = uniformLengthSized(&uniforms, offsets[1]); i++;

            out[i] = workgroupLengthUnsized(&wg1, offsets[0], bufferLength(&wg1)); i++;
            out[i] = workgroupLengthUnsized(&wg1, offsets[1], offsets[2]); i++;
            out[i] = workgroupLengthUnsized(&wg1, offsets[3], offsets[3]); i++;
            out[i] = workgroupLengthUnsized(&wg2, offsets[0], bufferLength(&wg2)); i++;
            out[i] = workgroupLengthSmall(&wg2, offsets[1]); i++;
            out[i] = workgroupLengthSmall(&wg2, offsets[2]); i++;
            out[i] = workgroupLengthUnsized(&wg3, offsets[0], bufferLength(&wg3)); i++;
            out[i] = workgroupLengthOverride(&wg3, offsets[1]); i++;
            out[i] = workgroupLengthOverride(&wg3, offsets[2]); i++;

            out[i] = storageLengthUnsized(&in, offsets[0], bufferLength(&in)); i++;
            out[i] = storageLengthUnsized(&in, offsets[0], offsets[4]); i++;
            out[i] = storageLengthUnsized(&in, offsets[1], bufferLength(&in) - 8); i++;
            out[i] = storageLengthSmall(&in, 0); i++;
          }
        )";

        std::initializer_list<uint32_t> uniforms = {0, 8, 16, 32, 64};
        std::vector<uint32_t> uniformData{uniforms};

        uint32_t uniformSize = static_cast<uint32_t>(uniforms.size());
        std::vector<uint32_t> expected = {
            // Uniform sizes
            uniformSize,
            uniformSize - uniformData[1] / 4,
            uniformSize - uniformData[2] / 4,
            uniformSize - 1,
            uniformSize - 1 - uniformData[1] / 4,
            // Workgroup sizes
            64 / 4,
            uniformData[2] / 4,
            uniformData[3] / 4,
            128 / 4,
            32 / 4,
            32 / 4,
            len / 4,
            32 / 4,
            32 / 4,
            // Storage sizes
            (len - 8) / 4,
            (uniformData[4] - 8) / 4,
            (len - 8 - uniformData[1]) / 4,
            (48 - 8) / 4,
        };

        BaseLengthTest(wgsl, len, uniforms, expected);
    }

    void BaseViewTest(const std::string& wgsl) {
        wgpu::ShaderModule module = utils::CreateShaderModule(device, wgsl);

        std::initializer_list<uint32_t> uniforms = {0, 1, 2,  3,  4,  5,  6,  7,
                                                    8, 9, 10, 11, 12, 13, 14, 15};
        std::initializer_list<uint32_t> inputs = {
            1, 2, 3, 4, 0x3f800000, 0x40000000, 0x40400000, 0x40800000, 7, 8, 9, 10, 11, 12, 13, 14,
        };
        const std::vector<uint32_t> expected = {
            1, 2, 3, 4, 0x3f800000, 0x40000000, 0x40400000, 0x40800000, 7, 8, 9, 10, 11, 12, 13, 14,
        };

        wgpu::ComputePipelineDescriptor desc;
        desc.compute.module = module;
        wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&desc);

        wgpu::Buffer src = utils::CreateBufferFromData<uint32_t>(
            device, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc, inputs);

        wgpu::BufferDescriptor outDesc;
        outDesc.size = sizeof(uint32_t) * expected.size();
        outDesc.usage =
            wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
        wgpu::Buffer out = device.CreateBuffer(&outDesc);

        wgpu::Buffer uni = utils::CreateBufferFromData<uint32_t>(
            device, wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopySrc, uniforms);

        wgpu::BindGroup bg =
            utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
                                 {
                                     {0, src, 0, inputs.size() * sizeof(uint32_t)},
                                     {1, out, 0, outDesc.size},
                                     {2, uni, 0, uniforms.size() * sizeof(uint32_t)},
                                 });

        wgpu::CommandBuffer commands;
        {
            wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
            wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
            pass.SetPipeline(pipeline);
            pass.SetBindGroup(0, bg);
            pass.DispatchWorkgroups(1, 1, 1);
            pass.End();
            commands = encoder.Finish();
        }
        queue.Submit(1, &commands);

        EXPECT_BUFFER_U32_RANGE_EQ(&expected[0], out, 0, expected.size());
    }

    void TestBufferViewBasics() {
        const std::string wgsl = R"(
          @group(0) @binding(0) var<storage> in : buffer;
          @group(0) @binding(1) var<storage, read_write> out : buffer;
          @group(0) @binding(2) var<uniform> uniforms : buffer<64>;

          struct S {
            a : u32,
            b : vec4u,
          }

          @compute @workgroup_size(1)
          fn main() {
            let uniform_v1 = bufferView<array<u32>>(&uniforms, 0);

            let in1 = bufferView<array<u32, 4>>(&in, uniform_v1[0]);
            let out1 = bufferView<vec4u>(&out, uniform_v1[0]);
            out1.x = in1[0];
            out1.y = in1[1];
            out1.z = in1[2];
            out1.w = in1[3];

            let in2x = bufferView<f32>(&in, uniform_v1[4] * 4);
            let in2y = bufferView<f32>(&in, uniform_v1[5] * 4);
            let in2z = bufferView<f32>(&in, uniform_v1[6] * 4);
            let in2w = bufferView<f32>(&in, uniform_v1[7] * 4);
            let out2 = bufferView<array<f32, 4>>(&out, 16u);
            out2[0] = *in2x;
            out2[1] = *in2y;
            out2[2] = *in2z;
            out2[3] = *in2w;

            let in3 = bufferView<S>(&in, uniform_v1[8] * 4);
            let in3_sub = bufferView<array<u32, 3>>(&in, uniform_v1[9] * 4);
            let out3 = bufferView<S>(&out, uniform_v1[8] * 4);
            let out3_sub = bufferView<array<u32, 3>>(&out, uniform_v1[9] * 4);
            let in3_ld = *in3;
            *out3_sub = *in3_sub;
            *out3 = in3_ld;
          }
        )";

        BaseViewTest(wgsl);
    }

    void TestBufferArrayViewBasics() {
        const std::string wgsl = R"(
          @group(0) @binding(0) var<storage> in : buffer;
          @group(0) @binding(1) var<storage, read_write> out : buffer;
          @group(0) @binding(2) var<uniform> uniforms : buffer<64>;

          struct S {
            a : u32,
            b : vec4u,
          }

          @compute @workgroup_size(1)
          fn main() {
            let uniform_v1 = bufferArrayView<array<u32>>(&uniforms, 0, 16 * 4);

            let in1 = bufferArrayView<array<u32>>(&in, uniform_v1[0], uniform_v1[5] * 4);
            let out1 = bufferArrayView<array<vec4u>>(&out, uniform_v1[0], 16);
            out1[0].x = in1[0];
            out1[0].y = in1[1];
            out1[0].z = in1[2];
            out1[0].w = in1[3];

            let in2x = bufferArrayView<array<f32>>(&in, uniform_v1[4] * 4, 8);
            let in2y = bufferArrayView<array<f32>>(&in, uniform_v1[5] * 4, 8);
            let in2z = bufferArrayView<array<f32>>(&in, uniform_v1[6] * 4, 8);
            let in2w = bufferArrayView<array<f32>>(&in, uniform_v1[7] * 4, 8);
            let out2 = bufferArrayView<array<f32>>(&out, 16u, 16u);
            out2[0] = in2x[0];
            out2[1] = in2y[0];
            out2[2] = in2z[0];
            out2[3] = in2w[0];

            let in3 = bufferArrayView<array<S>>(&in, uniform_v1[8] * 4, 32);
            let in3_sub = bufferArrayView<array<array<u32, 3>>>(&in, uniform_v1[9] * 4, 12);
            let out3 = bufferArrayView<array<S>>(&out, uniform_v1[8] * 4, 32);
            let out3_sub = bufferArrayView<array<array<u32, 3>>>(&out, uniform_v1[9] * 4, 12);
            let in3_ld = in3[0];
            out3_sub[0] = in3_sub[0];
            out3[0] = in3_ld;
          }
        )";

        BaseViewTest(wgsl);
    }
};

TEST_P(BufferViewTest, BufferViewArrayLength) {
    DAWN_SUPPRESS_TEST_IF(IsD3D11());
    DAWN_SUPPRESS_TEST_IF(IsD3D12());

    TestBufferViewArrayLength(128);
    TestBufferViewArrayLength(256);
    TestBufferViewArrayLength(512);
    TestBufferViewArrayLength(1024);
    TestBufferViewArrayLength(64);
}

TEST_P(BufferViewTest, BufferArrayViewArrayLength) {
    DAWN_SUPPRESS_TEST_IF(IsD3D11());
    DAWN_SUPPRESS_TEST_IF(IsD3D12());

    TestBufferArrayViewArrayLength(128);
    TestBufferArrayViewArrayLength(256);
    TestBufferArrayViewArrayLength(512);
    TestBufferArrayViewArrayLength(1024);
    TestBufferArrayViewArrayLength(64);
}

TEST_P(BufferViewTest, ArrayLengthFunctions) {
    DAWN_SUPPRESS_TEST_IF(IsD3D11());
    DAWN_SUPPRESS_TEST_IF(IsD3D12());

    TestArrayLengthFunctions(128);
    TestArrayLengthFunctions(256);
    TestArrayLengthFunctions(512);
    TestArrayLengthFunctions(1024);
    TestArrayLengthFunctions(64);
}

TEST_P(BufferViewTest, BufferViewBasics) {
    DAWN_SUPPRESS_TEST_IF(IsD3D11());
    DAWN_SUPPRESS_TEST_IF(IsD3D12());

    TestBufferViewBasics();
}

TEST_P(BufferViewTest, BufferArrayViewBasics) {
    DAWN_SUPPRESS_TEST_IF(IsD3D11());
    DAWN_SUPPRESS_TEST_IF(IsD3D12());

    TestBufferArrayViewBasics();
}

DAWN_INSTANTIATE_TEST(BufferViewTest,
                      D3D11Backend(),
                      D3D12Backend(),
                      MetalBackend(),
                      OpenGLBackend(),
                      OpenGLESBackend(),
                      VulkanBackend(),
                      WebGPUBackend());
}  // namespace
}  // namespace dawn