From fbda46da44a040ca00b178d0c74eb695d6364b4b Mon Sep 17 00:00:00 2001 From: "Yan, Shaobo" Date: Tue, 2 Mar 2021 03:00:27 +0000 Subject: [PATCH] Refactor CopyTextureForBrowserTests to use compute shader CopyTextureForBrowserTests will cover incoming color format conversion cases. While it is OK for |unorm| formats to compare their pixel values on the CPU side, we cannot do such comparisons for the |float| formats because we may meet the precision issues when comparing a value generated at the CPU side to the one from the GPU side. Refactor this test suites by using compute shader and do bit-by-bit comparison from source texture and destination texture. BUG=dawn:465 Change-Id: I979fcf1a1d96bbe9f8a4cf2f1a305d488e88b257 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/36140 Commit-Queue: Shaobo Yan Reviewed-by: Austin Eng --- .../end2end/CopyTextureForBrowserTests.cpp | 228 +++++++++++------- 1 file changed, 141 insertions(+), 87 deletions(-) diff --git a/src/tests/end2end/CopyTextureForBrowserTests.cpp b/src/tests/end2end/CopyTextureForBrowserTests.cpp index cca92c29ed..e37a80bb15 100644 --- a/src/tests/end2end/CopyTextureForBrowserTests.cpp +++ b/src/tests/end2end/CopyTextureForBrowserTests.cpp @@ -16,6 +16,7 @@ #include "common/Constants.h" #include "common/Math.h" +#include "utils/ComboRenderPipelineDescriptor.h" #include "utils/TestUtils.h" #include "utils/TextureFormatUtils.h" #include "utils/WGPUHelpers.h" @@ -23,14 +24,17 @@ class CopyTextureForBrowserTests : public DawnTest { protected: static constexpr wgpu::TextureFormat kTextureFormat = wgpu::TextureFormat::RGBA8Unorm; + static constexpr uint64_t kDefaultTextureWidth = 4; + static constexpr uint64_t kDefaultTextureHeight = 4; struct TextureSpec { - wgpu::Origin3D copyOrigin; - wgpu::Extent3D textureSize; - uint32_t level; + wgpu::Origin3D copyOrigin = {}; + wgpu::Extent3D textureSize = {kDefaultTextureWidth, kDefaultTextureHeight}; + uint32_t level = 0; + wgpu::TextureFormat format = kTextureFormat; }; - static std::vector GetExpectedTextureData(const utils::TextureDataCopyLayout& layout) { + static std::vector GetSourceTextureData(const utils::TextureDataCopyLayout& layout) { std::vector textureData(layout.texelBlockCount); for (uint32_t layer = 0; layer < layout.mipSize.depth; ++layer) { const uint32_t sliceOffset = layout.texelBlocksPerImage * layer; @@ -48,43 +52,87 @@ class CopyTextureForBrowserTests : public DawnTest { return textureData; } - static void PackTextureData(const RGBA8* srcData, - uint32_t width, - uint32_t height, - uint32_t srcTexelsPerRow, - RGBA8* dstData, - uint32_t dstTexelsPerRow, - const wgpu::CopyTextureForBrowserOptions* options) { - bool isFlipY = options != nullptr && options->flipY; - for (uint32_t y = 0; y < height; ++y) { - for (uint32_t x = 0; x < width; ++x) { - uint32_t srcYIndex = - isFlipY ? (height - y - 1) * srcTexelsPerRow : y * srcTexelsPerRow; - uint32_t src = x + srcYIndex; - uint32_t dst = x + y * dstTexelsPerRow; - dstData[dst] = srcData[src]; + void SetUp() override { + DawnTest::SetUp(); + + testPipeline = MakeTestPipeline(); + + uint32_t uniformBufferData[] = { + 0, // copy have flipY option + }; + + wgpu::BufferDescriptor uniformBufferDesc = {}; + uniformBufferDesc.usage = wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Uniform; + uniformBufferDesc.size = sizeof(uniformBufferData); + uniformBuffer = device.CreateBuffer(&uniformBufferDesc); + } + + // Do the bit-by-bit comparison between the source and destination texture with GPU (compute + // shader) instead of CPU after executing CopyTextureForBrowser() to avoid the errors caused by + // comparing a value generated on CPU to the one generated on GPU. + wgpu::ComputePipeline MakeTestPipeline() { + wgpu::ShaderModule csModule = utils::CreateShaderModuleFromWGSL(device, R"( + [[block]] struct Uniforms { + [[offset(0)]] dstTextureFlipY : u32; + }; + [[block]] struct OutputBuf { + [[offset(0)]] result : [[stride(4)]] array; + }; + [[group(0), binding(0)]] var src : texture_2d; + [[group(0), binding(1)]] var dst : texture_2d; + [[group(0), binding(2)]] var output : OutputBuf; + [[group(0), binding(3)]] var uniforms : Uniforms; + [[builtin(global_invocation_id)]] var GlobalInvocationID : vec3; + [[stage(compute), workgroup_size(1, 1, 1)]] + fn main() -> void { + // Current CopyTextureForBrowser only support full copy now. + // TODO(dawn:465): Refactor this after CopyTextureForBrowser + // support sub-rect copy. + var size : vec2 = textureDimensions(src); + var dstTexCoord : vec2 = vec2(GlobalInvocationID.xy); + var srcTexCoord : vec2 = dstTexCoord; + if (uniforms.dstTextureFlipY == 1) { + srcTexCoord.y = size.y - dstTexCoord.y - 1; + } + + var srcColor : vec4 = textureLoad(src, srcTexCoord); + var dstColor : vec4 = textureLoad(dst, dstTexCoord); + var success : bool = all(srcColor == dstColor); + + var outputIndex : u32 = GlobalInvocationID.y * size.x + GlobalInvocationID.x; + if (success) { + output.result[outputIndex] = u32(1); + } else { + output.result[outputIndex] = u32(0); + } } - } + )"); + + wgpu::ComputePipelineDescriptor csDesc; + csDesc.computeStage.module = csModule; + csDesc.computeStage.entryPoint = "main"; + + return device.CreateComputePipeline(&csDesc); } void DoTest(const TextureSpec& srcSpec, const TextureSpec& dstSpec, - const wgpu::Extent3D& copySize, - const wgpu::CopyTextureForBrowserOptions* options) { + const wgpu::Extent3D& copySize = {kDefaultTextureWidth, kDefaultTextureHeight}, + const wgpu::CopyTextureForBrowserOptions options = {}) { wgpu::TextureDescriptor srcDescriptor; srcDescriptor.size = srcSpec.textureSize; - srcDescriptor.format = kTextureFormat; + srcDescriptor.format = srcSpec.format; srcDescriptor.mipLevelCount = srcSpec.level + 1; - srcDescriptor.usage = wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | - wgpu::TextureUsage::Sampled | wgpu::TextureUsage::OutputAttachment; + srcDescriptor.usage = + wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | wgpu::TextureUsage::Sampled; wgpu::Texture srcTexture = device.CreateTexture(&srcDescriptor); wgpu::Texture dstTexture; wgpu::TextureDescriptor dstDescriptor; dstDescriptor.size = dstSpec.textureSize; - dstDescriptor.format = kTextureFormat; + dstDescriptor.format = dstSpec.format; dstDescriptor.mipLevelCount = dstSpec.level + 1; - dstDescriptor.usage = wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | + dstDescriptor.usage = wgpu::TextureUsage::CopyDst | wgpu::TextureUsage::Sampled | wgpu::TextureUsage::OutputAttachment; dstTexture = device.CreateTexture(&dstDescriptor); @@ -96,7 +144,7 @@ class CopyTextureForBrowserTests : public DawnTest { {srcSpec.textureSize.width, srcSpec.textureSize.height, copySize.depth}, srcSpec.level); - const std::vector textureArrayCopyData = GetExpectedTextureData(copyLayout); + const std::vector textureArrayCopyData = GetSourceTextureData(copyLayout); wgpu::TextureCopyView textureCopyView = utils::CreateTextureCopyView(srcTexture, srcSpec.level, {0, 0, srcSpec.copyOrigin.z}); @@ -109,7 +157,6 @@ class CopyTextureForBrowserTests : public DawnTest { textureArrayCopyData.size() * sizeof(RGBA8), &textureDataLayout, ©Layout.mipSize); - const wgpu::Extent3D copySizePerSlice = {copySize.width, copySize.height, 1}; // Perform the texture to texture copy wgpu::TextureCopyView srcTextureCopyView = utils::CreateTextureCopyView(srcTexture, srcSpec.level, srcSpec.copyOrigin); @@ -119,41 +166,64 @@ class CopyTextureForBrowserTests : public DawnTest { wgpu::CommandBuffer commands = encoder.Finish(); queue.Submit(1, &commands); - // Perform a copy here for testing. device.GetQueue().CopyTextureForBrowser(&srcTextureCopyView, &dstTextureCopyView, ©Size, - options); + &options); - // Texels in single slice. - const uint32_t texelCountInCopyRegion = utils::GetTexelCountInCopyRegion( - copyLayout.bytesPerRow, copyLayout.bytesPerImage / copyLayout.bytesPerRow, - copySizePerSlice, kTextureFormat); - std::vector expected(texelCountInCopyRegion); - for (uint32_t slice = 0; slice < copySize.depth; ++slice) { - std::fill(expected.begin(), expected.end(), RGBA8()); - const uint32_t texelIndexOffset = copyLayout.texelBlocksPerImage * slice; - const uint32_t expectedTexelArrayDataStartIndex = - texelIndexOffset + - (srcSpec.copyOrigin.x + srcSpec.copyOrigin.y * copyLayout.texelBlocksPerRow); - PackTextureData(&textureArrayCopyData[expectedTexelArrayDataStartIndex], copySize.width, - copySize.height, copyLayout.texelBlocksPerRow, expected.data(), - copySize.width, options); + // Update uniform buffer based on test config + uint32_t uniformBufferData[] = { + options.flipY, // copy have flipY option + }; - EXPECT_TEXTURE_RGBA8_EQ(expected.data(), dstTexture, dstSpec.copyOrigin.x, - dstSpec.copyOrigin.y, copySize.width, copySize.height, - dstSpec.level, dstSpec.copyOrigin.z + slice) - << "Texture to Texture copy failed copying region [(" << srcSpec.copyOrigin.x - << ", " << srcSpec.copyOrigin.y << "), (" << srcSpec.copyOrigin.x + copySize.width - << ", " << srcSpec.copyOrigin.y + copySize.height << ")) from " - << srcSpec.textureSize.width << " x " << srcSpec.textureSize.height - << " texture at mip level " << srcSpec.level << " layer " - << srcSpec.copyOrigin.z + slice << " to [(" << dstSpec.copyOrigin.x << ", " - << dstSpec.copyOrigin.y << "), (" << dstSpec.copyOrigin.x + copySize.width << ", " - << dstSpec.copyOrigin.y + copySize.height << ")) region of " - << dstSpec.textureSize.width << " x " << dstSpec.textureSize.height - << " texture at mip level " << dstSpec.level << " layer " - << dstSpec.copyOrigin.z + slice << std::endl; + device.GetQueue().WriteBuffer(uniformBuffer, 0, uniformBufferData, + sizeof(uniformBufferData)); + + // Create output buffer to store result + wgpu::BufferDescriptor outputDesc; + outputDesc.size = copySize.width * copySize.height * sizeof(uint32_t); + outputDesc.usage = + wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst; + wgpu::Buffer outputBuffer = device.CreateBuffer(&outputDesc); + + // Create texture views for test. + wgpu::TextureViewDescriptor srcTextureViewDesc = {}; + srcTextureViewDesc.baseMipLevel = srcSpec.level; + wgpu::TextureView srcTextureView = srcTexture.CreateView(&srcTextureViewDesc); + + wgpu::TextureViewDescriptor dstTextureViewDesc = {}; + dstTextureViewDesc.baseMipLevel = dstSpec.level; + wgpu::TextureView dstTextureView = dstTexture.CreateView(&dstTextureViewDesc); + + // Create bind group based on the config. + wgpu::BindGroup bindGroup = utils::MakeBindGroup( + device, testPipeline.GetBindGroupLayout(0), + {{0, srcTextureView}, + {1, dstTextureView}, + {2, outputBuffer, 0, copySize.width * copySize.height * sizeof(uint32_t)}, + {3, uniformBuffer, 0, sizeof(uniformBufferData)}}); + + // Start a pipeline to check pixel value in bit form. + wgpu::CommandEncoder testEncoder = device.CreateCommandEncoder(); + + wgpu::CommandBuffer testCommands; + { + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(testPipeline); + pass.SetBindGroup(0, bindGroup); + pass.Dispatch(copySize.width, copySize.height); + pass.EndPass(); + + testCommands = encoder.Finish(); } + queue.Submit(1, &testCommands); + + std::vector expectResult(copySize.width * copySize.height, 1); + EXPECT_BUFFER_U32_RANGE_EQ(expectResult.data(), outputBuffer, 0, + copySize.width * copySize.height); } + + wgpu::Buffer uniformBuffer; // Uniform buffer to store dst texture meta info. + wgpu::ComputePipeline testPipeline; }; // Verify CopyTextureForBrowserTests works with internal pipeline. @@ -166,12 +236,9 @@ TEST_P(CopyTextureForBrowserTests, PassthroughCopy) { constexpr uint32_t kHeight = 1; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; - wgpu::CopyTextureForBrowserOptions options = {}; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}); } TEST_P(CopyTextureForBrowserTests, VerifyCopyOnXDirection) { @@ -182,12 +249,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyOnXDirection) { constexpr uint32_t kHeight = 1; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; - wgpu::CopyTextureForBrowserOptions options = {}; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}); } TEST_P(CopyTextureForBrowserTests, VerifyCopyOnYDirection) { @@ -198,12 +262,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyOnYDirection) { constexpr uint32_t kHeight = 1000; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; - wgpu::CopyTextureForBrowserOptions options = {}; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}); } TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) { @@ -214,12 +275,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) { constexpr uint32_t kHeight = 999; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; - wgpu::CopyTextureForBrowserOptions options = {}; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}); } TEST_P(CopyTextureForBrowserTests, VerifyFlipY) { @@ -230,13 +288,11 @@ TEST_P(CopyTextureForBrowserTests, VerifyFlipY) { constexpr uint32_t kHeight = 1001; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; wgpu::CopyTextureForBrowserOptions options = {}; options.flipY = true; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}, options); } TEST_P(CopyTextureForBrowserTests, VerifyFlipYInSlimTexture) { @@ -247,13 +303,11 @@ TEST_P(CopyTextureForBrowserTests, VerifyFlipYInSlimTexture) { constexpr uint32_t kHeight = 1001; TextureSpec textureSpec; - textureSpec.copyOrigin = {0, 0, 0}; - textureSpec.level = 0; - textureSpec.textureSize = {kWidth, kHeight, 1}; + textureSpec.textureSize = {kWidth, kHeight}; wgpu::CopyTextureForBrowserOptions options = {}; options.flipY = true; - DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); + DoTest(textureSpec, textureSpec, {kWidth, kHeight}, options); } DAWN_INSTANTIATE_TEST(CopyTextureForBrowserTests,