[YUV AHB] Add ExternalTextureTests using multiple textures Similar tests will be added for the YCbCr texture sampling to check that the JITing works correctly. It is useful to compare them between YCbCr and non-YCbCr cases, especially because a future CL makes the RGBA textures take mostly the same path as the YCbCr ones. Bug: 468988322 Change-Id: I275514efb993ea6a1c7d94b1194c60f943f554da Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/299375 Commit-Queue: Corentin Wallez <cwallez@chromium.org> Reviewed-by: Brandon Jones <bajones@chromium.org> Reviewed-by: Kyle Charbonneau <kylechar@google.com>
diff --git a/src/dawn/tests/DawnTest.h b/src/dawn/tests/DawnTest.h index 8a93fcc..e7dccf7 100644 --- a/src/dawn/tests/DawnTest.h +++ b/src/dawn/tests/DawnTest.h
@@ -93,6 +93,14 @@ EXPECT_BUFFER(buffer, offset, sizeof(uint32_t) * (count), \ new ::dawn::detail::ExpectEq<uint32_t>(expected, count)) +#define EXPECT_BUFFER_RGBA8_EQ(expected, buffer, offset) \ + EXPECT_BUFFER(buffer, offset, sizeof(utils::RGBA8), \ + new ::dawn::detail::ExpectEq<utils::RGBA8>(expected)) + +#define EXPECT_BUFFER_RGBA8_RANGE_EQ(expected, buffer, offset, count) \ + EXPECT_BUFFER(buffer, offset, sizeof(utils::RGBA8) * (count), \ + new ::dawn::detail::ExpectEq<utils::RGBA8>(expected, count)) + #define EXPECT_BUFFER_U64_EQ(expected, buffer, offset) \ EXPECT_BUFFER(buffer, offset, sizeof(uint64_t), \ new ::dawn::detail::ExpectEq<uint64_t>(expected))
diff --git a/src/dawn/tests/end2end/ExternalTextureTests.cpp b/src/dawn/tests/end2end/ExternalTextureTests.cpp index de72c1f..85793d2 100644 --- a/src/dawn/tests/end2end/ExternalTextureTests.cpp +++ b/src/dawn/tests/end2end/ExternalTextureTests.cpp
@@ -200,6 +200,39 @@ this->queue.Submit(1, &commands); } + wgpu::Texture MakeTestTexture(wgpu::TextureFormat format, + uint32_t width, + uint32_t height, + wgpu::Color color) { + wgpu::TextureDescriptor tDesc = { + .usage = wgpu::TextureUsage::RenderAttachment | wgpu::TextureUsage::TextureBinding, + .size = {width, height}, + .format = format, + }; + wgpu::Texture texture = this->device.CreateTexture(&tDesc); + + // Use a render pass clear to set the value for the texture so that it works whatever the + // format. + wgpu::RenderPassColorAttachment attachment = { + .view = texture.CreateView(), + .loadOp = wgpu::LoadOp::Clear, + .storeOp = wgpu::StoreOp::Store, + .clearValue = color, + }; + wgpu::RenderPassDescriptor rpDesc = { + .colorAttachmentCount = 1, + .colorAttachments = &attachment, + }; + + wgpu::CommandEncoder encoder = this->device.CreateCommandEncoder(); + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&rpDesc); + pass.End(); + wgpu::CommandBuffer commands = encoder.Finish(); + this->queue.Submit(1, &commands); + + return texture; + } + static constexpr uint32_t kWidth = 4; static constexpr uint32_t kHeight = 4; static constexpr wgpu::TextureFormat kFormat = wgpu::TextureFormat::RGBA8Unorm; @@ -1590,6 +1623,168 @@ ASSERT_NE(pipeline.Get(), nullptr); } +// Tests using multiple textures in the same pipeline and checks that sampling them returns the +// expected result. It rolls the RGBA and multiplanar external textures to check that the Vulkan +// pipeline specialization works correctly (and doesn't reuse when it shouldn't). +// +// Case with all in the same bindgroup layout. +TEST_P(ExternalTextureTests, SampleDifferentKindsSameBindGroup) { + DAWN_SUPPRESS_TEST_IF(IsWARP()); + + // Create our three test external of different kinds as well as the expected data. + std::vector<wgpu::ExternalTexture> externalTextures; + std::vector<utils::RGBA8> colors; + + // The RGBA external textures. + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::RGBA8Unorm, 1, 1, + {1.0 / 255.0, 2.0 / 255.0, 3.0 / 255.0, 4.0 / 255.0}))); + colors.push_back(utils::RGBA8(1, 2, 3, 4)); + + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::RGBA8Unorm, 1, 1, + {2.0 / 255.0, 4.0 / 255.0, 6.0 / 255.0, 8.0 / 255.0}))); + colors.push_back(utils::RGBA8(2, 4, 6, 8)); + + // The multiplanar external texture. + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::R8Unorm, 2, 2, {10.0 / 255.0, 0, 0, 0}), + MakeTestTexture(wgpu::TextureFormat::RG8Unorm, 1, 1, {20.0 / 255.0, 30.0 / 255.0, 0, 0}))); + colors.push_back(utils::RGBA8(10, 20, 30, 255)); + + // Create the pipeline that copies from the three external texture to a buffer. + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = utils::CreateShaderModule(device, R"( + @group(0) @binding(0) var s : sampler; + @group(0) @binding(1) var<storage, read_write> results : array<u32, 3>; + + @group(0) @binding(2) var t0 : texture_external; + @group(0) @binding(3) var t1 : texture_external; + @group(0) @binding(4) var t2 : texture_external; + @compute @workgroup_size(1) fn main() { + results[0] = pack4x8unorm(textureSampleBaseClampToEdge(t0, s, vec2(0))); + results[1] = pack4x8unorm(textureSampleBaseClampToEdge(t1, s, vec2(0))); + results[2] = pack4x8unorm(textureSampleBaseClampToEdge(t2, s, vec2(0))); + } + )"); + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc); + + wgpu::BufferDescriptor resultDesc = { + .usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc, + .size = 3 * sizeof(uint32_t), + }; + wgpu::Buffer resultBuffer = device.CreateBuffer(&resultDesc); + + // Run the pipeline, rolling the external textures in the bindgroup such that a new + // specialization of the shader should be created each time. + for (size_t roll = 0; roll < externalTextures.size(); roll++) { + wgpu::BindGroup bg = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + { + {0, device.CreateSampler()}, + {1, resultBuffer}, + {2, externalTextures[(0 + roll) % 3]}, + {3, externalTextures[(1 + roll) % 3]}, + {4, externalTextures[(2 + roll) % 3]}, + }); + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetBindGroup(0, bg); + pass.SetPipeline(pipeline); + pass.DispatchWorkgroups(1); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_BUFFER_RGBA8_EQ(colors[(0 + roll) % 3], resultBuffer, 0); + EXPECT_BUFFER_RGBA8_EQ(colors[(1 + roll) % 3], resultBuffer, 4); + EXPECT_BUFFER_RGBA8_EQ(colors[(2 + roll) % 3], resultBuffer, 8); + } +} + +// Case with all in different bind group layouts. +TEST_P(ExternalTextureTests, SampleDifferentKindsDifferentBindGroups) { + DAWN_SUPPRESS_TEST_IF(IsWARP()); + + // Create our three test external of different kinds as well as the expected data. + std::vector<wgpu::ExternalTexture> externalTextures; + std::vector<utils::RGBA8> colors; + + // The RGBA external textures. + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::RGBA8Unorm, 1, 1, + {1.0 / 255.0, 2.0 / 255.0, 3.0 / 255.0, 4.0 / 255.0}))); + colors.push_back(utils::RGBA8(1, 2, 3, 4)); + + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::RGBA8Unorm, 1, 1, + {2.0 / 255.0, 4.0 / 255.0, 6.0 / 255.0, 8.0 / 255.0}))); + colors.push_back(utils::RGBA8(2, 4, 6, 8)); + + // The multiplanar external texture. + externalTextures.push_back(utils::MakePassthroughExternalTexture( + device, MakeTestTexture(wgpu::TextureFormat::R8Unorm, 2, 2, {10.0 / 255.0, 0, 0, 0}), + MakeTestTexture(wgpu::TextureFormat::RG8Unorm, 1, 1, {20.0 / 255.0, 30.0 / 255.0, 0, 0}))); + colors.push_back(utils::RGBA8(10, 20, 30, 255)); + + // Create the pipeline that copies from the three external texture to a buffer. + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = utils::CreateShaderModule(device, R"( + @group(0) @binding(0) var s : sampler; + @group(0) @binding(1) var<storage, read_write> results : array<u32, 3>; + + @group(1) @binding(0) var t0 : texture_external; + @group(2) @binding(0) var t1 : texture_external; + @group(3) @binding(0) var t2 : texture_external; + @compute @workgroup_size(1) fn main() { + results[0] = pack4x8unorm(textureSampleBaseClampToEdge(t0, s, vec2(0))); + results[1] = pack4x8unorm(textureSampleBaseClampToEdge(t1, s, vec2(0))); + results[2] = pack4x8unorm(textureSampleBaseClampToEdge(t2, s, vec2(0))); + } + )"); + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc); + + wgpu::BufferDescriptor resultDesc = { + .usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc, + .size = 3 * sizeof(uint32_t), + }; + wgpu::Buffer resultBuffer = device.CreateBuffer(&resultDesc); + + // Run the pipeline, rolling the external textures in the bindgroup such that a new + // specialization of the shader should be created each time. + for (size_t roll = 0; roll < externalTextures.size(); roll++) { + wgpu::BindGroup bg0 = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + { + {0, device.CreateSampler()}, + {1, resultBuffer}, + }); + + wgpu::BindGroup bg1 = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(1), + {{0, externalTextures[(0 + roll) % 3]}}); + wgpu::BindGroup bg2 = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(2), + {{0, externalTextures[(1 + roll) % 3]}}); + wgpu::BindGroup bg3 = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(3), + {{0, externalTextures[(2 + roll) % 3]}}); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetBindGroup(0, bg0); + pass.SetBindGroup(1, bg1); + pass.SetBindGroup(2, bg2); + pass.SetBindGroup(3, bg3); + pass.SetPipeline(pipeline); + pass.DispatchWorkgroups(1); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_BUFFER_RGBA8_EQ(colors[(0 + roll) % 3], resultBuffer, 0); + EXPECT_BUFFER_RGBA8_EQ(colors[(1 + roll) % 3], resultBuffer, 4); + EXPECT_BUFFER_RGBA8_EQ(colors[(2 + roll) % 3], resultBuffer, 8); + } +} + DAWN_INSTANTIATE_TEST(ExternalTextureTests, D3D11Backend(), D3D12Backend(),
diff --git a/src/dawn/utils/WGPUHelpers.cpp b/src/dawn/utils/WGPUHelpers.cpp index dea8722..6ecf6f2 100644 --- a/src/dawn/utils/WGPUHelpers.cpp +++ b/src/dawn/utils/WGPUHelpers.cpp
@@ -456,6 +456,26 @@ return info; } +#ifndef __EMSCRIPTEN__ +wgpu::ExternalTexture MakePassthroughExternalTexture(const wgpu::Device& device, + const wgpu::Texture& plane0, + const wgpu::Texture& plane1) { + utils::ColorSpaceConversionInfo noopConversion = utils::GetNoopColorSpaceConversionInfo(); + wgpu::ExternalTextureDescriptor etDesc = { + .plane0 = plane0.CreateView(), + .plane1 = plane1 ? plane1.CreateView() : nullptr, + .cropOrigin = {0, 0}, + .cropSize = {plane0.GetWidth(), plane0.GetHeight()}, + .apparentSize = {plane0.GetWidth(), plane0.GetHeight()}, + .yuvToRgbConversionMatrix = noopConversion.yuvToRgbConversionMatrix.data(), + .srcTransferFunctionParameters = noopConversion.srcTransferFunctionParameters.data(), + .dstTransferFunctionParameters = noopConversion.dstTransferFunctionParameters.data(), + .gamutConversionMatrix = noopConversion.gamutConversionMatrix.data(), + }; + return device.CreateExternalTexture(&etDesc); +} +#endif // __EMSCRIPTEN__ + bool BackendRequiresCompat(wgpu::BackendType backend) { switch (backend) { case wgpu::BackendType::D3D12:
diff --git a/src/dawn/utils/WGPUHelpers.h b/src/dawn/utils/WGPUHelpers.h index d23bc42..966d685 100644 --- a/src/dawn/utils/WGPUHelpers.h +++ b/src/dawn/utils/WGPUHelpers.h
@@ -224,6 +224,14 @@ ColorSpaceConversionInfo GetYUVBT709ToRGBSRGBColorSpaceConversionInfo(); ColorSpaceConversionInfo GetNoopColorSpaceConversionInfo(); +#ifndef __EMSCRIPTEN__ +// Make an external texture from one or two planes that doesn't perform any color-space conversion +// or YUV to RGB conversion. The planes are given as textures so that we can reflect their size. +wgpu::ExternalTexture MakePassthroughExternalTexture(const wgpu::Device& device, + const wgpu::Texture& plane0, + const wgpu::Texture& plane1 = {}); +#endif // __EMSCRIPTEN__ + bool BackendRequiresCompat(wgpu::BackendType backend); absl::flat_hash_set<wgpu::FeatureName> FeatureAndImplicitlyEnabled(wgpu::FeatureName featureName);