| // 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 |
