CopyTextureForBrowser: Support Subrect Copy

This CL support subrect copy. Use scale/offset to copy from subrect of
source texture and viewport for copy to subrect of dstTexture.

BUG=dawn:465

Change-Id: Ice43c0da15f6d9526912879e2e734f6570f2d673
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/46422
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Shaobo Yan <shaobo.yan@intel.com>

src/dawn_native/CopyTextureForBrowserHelper.cpp

@@ -49,10 +49,25 @@ namespace dawn_native {
             [[stage(vertex)]] fn main() -> void {
                 Position = vec4<f32>((texcoord[VertexIndex] * 2.0 - vec2<f32>(1.0, 1.0)), 0.0, 1.0);
 
+                // Y component of scale is calculated by the copySizeHeight / textureHeight. Only
+                // flipY case can get negative number.
+                var flipY : bool = uniforms.u_scale.y < 0.0;
+
                 // Texture coordinate takes top-left as origin point. We need to map the
                 // texture to triangle carefully.
-                v_texcoord = (texcoord[VertexIndex] * vec2<f32>(1.0, -1.0) + vec2<f32>(0.0, 1.0)) *
+                if (flipY) {
-                    uniforms.u_scale + uniforms.u_offset;
+                    // We need to get the mirror positions(mirrored based on y = 0.5) on flip cases.
+                    // Adopt transform to src texture and then mapping it to triangle coord which
+                    // do a +1 shift on Y dimension will help us got that mirror position perfectly.
+                    v_texcoord = (texcoord[VertexIndex] * uniforms.u_scale + uniforms.u_offset) *
+                                  vec2<f32>(1.0, -1.0) + vec2<f32>(0.0, 1.0);
+                } else {
+                    // For the normal case, we need to get the exact position.
+                    // So mapping texture to triangle firstly then adopt the transform.
+                    v_texcoord = (texcoord[VertexIndex] *
+                                  vec2<f32>(1.0, -1.0) + vec2<f32>(0.0, 1.0)) *
+                                  uniforms.u_scale + uniforms.u_offset;
+                }
             }
         )";
 
@@ -219,17 +234,6 @@ namespace dawn_native {
 
         DAWN_TRY(ValidateCopyTextureForBrowserOptions(options));
 
-        // TODO(shaobo.yan@intel.com): Support the simplest case for now that source and destination
-        // texture has the same size and do full texture blit. Will address sub texture blit in
-        // future and remove these validations.
-        if (source->origin.x != 0 || source->origin.y != 0 || source->origin.z != 0 ||
-            destination->origin.x != 0 || destination->origin.y != 0 ||
-            destination->origin.z != 0 || source->mipLevel != 0 || destination->mipLevel != 0 ||
-            source->texture->GetWidth() != destination->texture->GetWidth() ||
-            source->texture->GetHeight() != destination->texture->GetHeight()) {
-            return DAWN_VALIDATION_ERROR("Cannot support sub blit now.");
-        }
-
         return {};
     }
 
@@ -255,16 +259,22 @@ namespace dawn_native {
         bgDesc.entryCount = 3;
         bgDesc.entries = bindGroupEntries;
 
+        Extent3D srcTextureSize = source->texture->GetSize();
+
         // Prepare binding 0 resource: uniform buffer.
         float uniformData[] = {
-            1.0, 1.0,  // scale
+            copySize->width / static_cast<float>(srcTextureSize.width),
-            0.0, 0.0   // offset
+            copySize->height / static_cast<float>(srcTextureSize.height),  // scale
+            source->origin.x / static_cast<float>(srcTextureSize.width),
+            source->origin.y / static_cast<float>(srcTextureSize.height)  // offset
         };
 
-        // Handle flipY
+        // Handle flipY. FlipY here means we flip the source texture firstly and then
+        // do copy. This helps on the case which source texture is flipped and the copy
+        // need to unpack the flip.
         if (options && options->flipY) {
             uniformData[1] *= -1.0;
-            uniformData[3] += 1.0;
+            uniformData[3] += copySize->height / static_cast<float>(srcTextureSize.height);
         }
 
         BufferDescriptor uniformDesc = {};
@@ -336,6 +346,8 @@ namespace dawn_native {
         // the copy from src texture to dst texture with transformation.
         passEncoder->APISetPipeline(pipeline);
         passEncoder->APISetBindGroup(0, bindGroup.Get());
+        passEncoder->APISetViewport(destination->origin.x, destination->origin.y, copySize->width,
+                                    copySize->height, 0.0, 1.0);
         passEncoder->APIDraw(3);
         passEncoder->APIEndPass();
 

src/tests/end2end/CopyTextureForBrowserTests.cpp

@@ -34,6 +34,10 @@ namespace {
         wgpu::TextureFormat::RGBA32Float, wgpu::TextureFormat::RG8Unorm,
         wgpu::TextureFormat::RGBA16Float, wgpu::TextureFormat::RG16Float,
         wgpu::TextureFormat::RGB10A2Unorm};
+
+    static const wgpu::Origin3D kOrigins[] = {{1, 1}, {1, 2}, {2, 1}};
+
+    static const wgpu::Extent3D kCopySize[] = {{1, 1}, {2, 1}, {1, 2}, {2, 2}};
 }  // anonymous namespace
 
 class CopyTextureForBrowserTests : public DawnTest {
@@ -72,17 +76,33 @@ class CopyTextureForBrowserTests : public DawnTest {
         return sourceTextureData;
     }
 
-    static std::vector<RGBA8> GetSourceTextureData(const utils::TextureDataCopyLayout& layout) {
+    enum class TextureCopyRole {
+        SOURCE,
+        DEST,
+    };
+
+    // Source texture contains red pixels and dst texture contains green pixels at start.
+    static std::vector<RGBA8> GetTextureData(const utils::TextureDataCopyLayout& layout,
+                                             TextureCopyRole textureRole) {
         std::vector<RGBA8> textureData(layout.texelBlockCount);
         for (uint32_t layer = 0; layer < layout.mipSize.depthOrArrayLayers; ++layer) {
             const uint32_t sliceOffset = layout.texelBlocksPerImage * layer;
             for (uint32_t y = 0; y < layout.mipSize.height; ++y) {
                 const uint32_t rowOffset = layout.texelBlocksPerRow * y;
                 for (uint32_t x = 0; x < layout.mipSize.width; ++x) {
-                    textureData[sliceOffset + rowOffset + x] =
+                    // Source textures will have variable pixel data to cover cases like
-                        RGBA8(static_cast<uint8_t>((x + layer * x) % 256),
+                    // flipY.
-                              static_cast<uint8_t>((y + layer * y) % 256),
+                    if (textureRole == TextureCopyRole::SOURCE) {
-                              static_cast<uint8_t>(x % 256), static_cast<uint8_t>(x % 256));
+                        textureData[sliceOffset + rowOffset + x] =
+                            RGBA8(static_cast<uint8_t>((x + layer * x) % 256),
+                                  static_cast<uint8_t>((y + layer * y) % 256),
+                                  static_cast<uint8_t>(x % 256), static_cast<uint8_t>(x % 256));
+                    } else {  // Dst textures will have be init as `green` to ensure subrect
+                              // copy not cross bound.
+                        textureData[sliceOffset + rowOffset + x] =
+                            RGBA8(static_cast<uint8_t>(0), static_cast<uint8_t>(255),
+                                  static_cast<uint8_t>(0), static_cast<uint8_t>(255));
+                    }
                 }
             }
         }
@@ -99,8 +119,11 @@ class CopyTextureForBrowserTests : public DawnTest {
         testPipeline = MakeTestPipeline();
 
         uint32_t uniformBufferData[] = {
-            0,  // copy have flipY option
+            0,     // copy have flipY option
-            4,  // channelCount
+            4,     // channelCount
+            0, 0,  // uvec2, subrect copy src origin
+            0, 0,  // uvec2, subrect copy dst origin
+            0, 0,  // uvec2, subrect copy size
         };
 
         wgpu::BufferDescriptor uniformBufferDesc = {};
@@ -116,7 +139,10 @@ class CopyTextureForBrowserTests : public DawnTest {
         wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
             [[block]] struct Uniforms {
                 dstTextureFlipY : u32;
-                channelCount : u32;
+                channelCount    : u32;
+                srcCopyOrigin   : vec2<u32>;
+                dstCopyOrigin   : vec2<u32>;
+                copySize        : vec2<u32>;
             };
             [[block]] struct OutputBuf {
                 result : array<u32>;
@@ -131,35 +157,50 @@ class CopyTextureForBrowserTests : public DawnTest {
                 return abs(value - expect) < 0.001;
             }
             [[stage(compute), workgroup_size(1, 1, 1)]] fn main() -> void {
-                // Current CopyTextureForBrowser only support full copy now.
+                var srcSize : vec2<i32> = textureDimensions(src);
-                // TODO(crbug.com/dawn/465): Refactor this after CopyTextureForBrowser
+                var dstSize : vec2<i32> = textureDimensions(dst);
-                // support sub-rect copy.
+                var dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
-                var size : vec2<i32> = textureDimensions(src);
+                var nonCoveredColor : vec4<f32> =
-                var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
+                    vec4<f32>(0.0, 1.0, 0.0, 1.0); // should be green
-                var srcTexCoord : vec2<i32> = dstTexCoord;
-                if (uniforms.dstTextureFlipY == 1u) {
-                    srcTexCoord.y = size.y - dstTexCoord.y - 1;
-                }
 
-                var srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
-                var dstColor : vec4<f32> = textureLoad(dst, dstTexCoord, 0);
                 var success : bool = true;
-
+                if (dstTexCoord.x < uniforms.dstCopyOrigin.x ||
-                // Not use loop and variable index format to workaround
+                    dstTexCoord.y < uniforms.dstCopyOrigin.y ||
-                // crbug.com/tint/638.
+                    dstTexCoord.x >= uniforms.dstCopyOrigin.x + uniforms.copySize.x ||
-                if (uniforms.channelCount == 2u) { // All have rg components.
+                    dstTexCoord.y >= uniforms.dstCopyOrigin.y + uniforms.copySize.y) {
                     success = success &&
-                              aboutEqual(dstColor.r, srcColor.r) &&
+                              all(textureLoad(dst, vec2<i32>(dstTexCoord), 0) == nonCoveredColor);
-                              aboutEqual(dstColor.g, srcColor.g);
                 } else {
-                    success = success &&
+                    // Calculate source texture coord.
-                              aboutEqual(dstColor.r, srcColor.r) &&
+                    var srcTexCoord : vec2<u32> = dstTexCoord - uniforms.dstCopyOrigin +
-                              aboutEqual(dstColor.g, srcColor.g) &&
+                                                  uniforms.srcCopyOrigin;
-                              aboutEqual(dstColor.b, srcColor.b) &&
+                    // Note that |flipY| equals flip src texture firstly and then do copy from src
-                              aboutEqual(dstColor.a, srcColor.a);
+                    // subrect to dst subrect. This helps on blink part to handle some input texture
-                }
+                    // which is flipped and need to unpack flip during the copy.
+                    // We need to calculate the expect y coord based on this rule.
+                    if (uniforms.dstTextureFlipY == 1u) {
+                        srcTexCoord.y = u32(srcSize.y) - srcTexCoord.y - 1u;
+                    }
 
-                var outputIndex : u32 = GlobalInvocationID.y * u32(size.x) + GlobalInvocationID.x;
+                    var srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
+                    var dstColor : vec4<f32> = textureLoad(dst, vec2<i32>(dstTexCoord), 0);
+
+                    // Not use loop and variable index format to workaround
+                    // crbug.com/tint/638.
+                    if (uniforms.channelCount == 2u) { // All have rg components.
+                        success = success &&
+                                  aboutEqual(dstColor.r, srcColor.r) &&
+                                  aboutEqual(dstColor.g, srcColor.g);
+                    } else {
+                        success = success &&
+                                  aboutEqual(dstColor.r, srcColor.r) &&
+                                  aboutEqual(dstColor.g, srcColor.g) &&
+                                  aboutEqual(dstColor.b, srcColor.b) &&
+                                  aboutEqual(dstColor.a, srcColor.a);
+                    }
+                }
+                var outputIndex : u32 = GlobalInvocationID.y * u32(dstSize.x) +
+                                        GlobalInvocationID.x;
                 if (success) {
                     output.result[outputIndex] = 1u;
                 } else {
@@ -199,6 +240,7 @@ class CopyTextureForBrowserTests : public DawnTest {
                 const wgpu::Extent3D& copySize = {kDefaultTextureWidth, kDefaultTextureHeight},
                 const wgpu::CopyTextureForBrowserOptions options = {},
                 bool useFixedTestValue = false) {
+        // Create and initialize src texture.
         wgpu::TextureDescriptor srcDescriptor;
         srcDescriptor.size = srcSpec.textureSize;
         srcDescriptor.format = srcSpec.format;
@@ -207,6 +249,41 @@ class CopyTextureForBrowserTests : public DawnTest {
             wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | wgpu::TextureUsage::Sampled;
         wgpu::Texture srcTexture = device.CreateTexture(&srcDescriptor);
 
+        const utils::TextureDataCopyLayout srcCopyLayout =
+            utils::GetTextureDataCopyLayoutForTexture2DAtLevel(
+                kTextureFormat,
+                {srcSpec.textureSize.width, srcSpec.textureSize.height,
+                 copySize.depthOrArrayLayers},
+                srcSpec.level);
+
+        std::vector<RGBA8> srcTextureArrayCopyData;
+        if (useFixedTestValue) {  // Use fixed value for color conversion tests.
+            srcTextureArrayCopyData = GetFixedSourceTextureData();
+        } else {  // For other tests, the input format is always kTextureFormat.
+
+            srcTextureArrayCopyData = GetTextureData(srcCopyLayout, TextureCopyRole::SOURCE);
+        }
+
+        wgpu::ImageCopyTexture srcImageTextureInit =
+            utils::CreateImageCopyTexture(srcTexture, srcSpec.level, {0, 0});
+
+        wgpu::TextureDataLayout srcTextureDataLayout;
+        srcTextureDataLayout.offset = 0;
+        srcTextureDataLayout.bytesPerRow = srcCopyLayout.bytesPerRow;
+        srcTextureDataLayout.rowsPerImage = srcCopyLayout.rowsPerImage;
+
+        device.GetQueue().WriteTexture(&srcImageTextureInit, srcTextureArrayCopyData.data(),
+                                       srcTextureArrayCopyData.size() * sizeof(RGBA8),
+                                       &srcTextureDataLayout, &srcCopyLayout.mipSize);
+
+        bool testSubRectCopy = srcSpec.copyOrigin.x > 0 || srcSpec.copyOrigin.y > 0 ||
+                               dstSpec.copyOrigin.x > 0 || dstSpec.copyOrigin.y > 0 ||
+                               srcSpec.textureSize.width > copySize.width ||
+                               srcSpec.textureSize.height > copySize.height ||
+                               dstSpec.textureSize.width > copySize.width ||
+                               dstSpec.textureSize.height > copySize.height;
+
+        // Create and init dst texture.
         wgpu::Texture dstTexture;
         wgpu::TextureDescriptor dstDescriptor;
         dstDescriptor.size = dstSpec.textureSize;
@@ -216,52 +293,57 @@ class CopyTextureForBrowserTests : public DawnTest {
                               wgpu::TextureUsage::OutputAttachment | wgpu::TextureUsage::CopySrc;
         dstTexture = device.CreateTexture(&dstDescriptor);
 
-        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        if (testSubRectCopy) {
+            // For subrect copy tests, dst texture use kTextureFormat always.
+            const utils::TextureDataCopyLayout dstCopyLayout =
+                utils::GetTextureDataCopyLayoutForTexture2DAtLevel(
+                    kTextureFormat,
+                    {dstSpec.textureSize.width, dstSpec.textureSize.height,
+                     copySize.depthOrArrayLayers},
+                    dstSpec.level);
 
-        const utils::TextureDataCopyLayout copyLayout =
+            const std::vector<RGBA8> dstTextureArrayCopyData =
-            utils::GetTextureDataCopyLayoutForTexture2DAtLevel(
+                GetTextureData(dstCopyLayout, TextureCopyRole::DEST);
-                kTextureFormat,
-                {srcSpec.textureSize.width, srcSpec.textureSize.height,
-                 copySize.depthOrArrayLayers},
-                srcSpec.level);
 
-        const std::vector<RGBA8> textureArrayCopyData =
+            wgpu::TextureDataLayout dstTextureDataLayout;
-            useFixedTestValue ? GetFixedSourceTextureData() : GetSourceTextureData(copyLayout);
+            dstTextureDataLayout.offset = 0;
-        wgpu::ImageCopyTexture imageCopyTexture =
+            dstTextureDataLayout.bytesPerRow = dstCopyLayout.bytesPerRow;
-            utils::CreateImageCopyTexture(srcTexture, srcSpec.level, {0, 0, srcSpec.copyOrigin.z});
+            dstTextureDataLayout.rowsPerImage = dstCopyLayout.rowsPerImage;
 
-        wgpu::TextureDataLayout textureDataLayout;
+            wgpu::ImageCopyTexture dstImageTextureInit =
-        textureDataLayout.offset = 0;
+                utils::CreateImageCopyTexture(dstTexture, dstSpec.level, {0, 0});
-        textureDataLayout.bytesPerRow = copyLayout.bytesPerRow;
-        textureDataLayout.rowsPerImage = copyLayout.rowsPerImage;
 
-        device.GetQueue().WriteTexture(&imageCopyTexture, textureArrayCopyData.data(),
+            device.GetQueue().WriteTexture(&dstImageTextureInit, dstTextureArrayCopyData.data(),
-                                       textureArrayCopyData.size() * sizeof(RGBA8),
+                                           dstTextureArrayCopyData.size() * sizeof(RGBA8),
-                                       &textureDataLayout, &copyLayout.mipSize);
+                                           &dstTextureDataLayout, &dstCopyLayout.mipSize);
+        }
 
         // Perform the texture to texture copy
         wgpu::ImageCopyTexture srcImageCopyTexture =
             utils::CreateImageCopyTexture(srcTexture, srcSpec.level, srcSpec.copyOrigin);
         wgpu::ImageCopyTexture dstImageCopyTexture =
             utils::CreateImageCopyTexture(dstTexture, dstSpec.level, dstSpec.copyOrigin);
-
-        wgpu::CommandBuffer commands = encoder.Finish();
-        queue.Submit(1, &commands);
-
         device.GetQueue().CopyTextureForBrowser(&srcImageCopyTexture, &dstImageCopyTexture,
                                                 &copySize, &options);
 
         // Update uniform buffer based on test config
         uint32_t uniformBufferData[] = {
-            options.flipY,                                    // copy have flipY option
+            options.flipY,                                   // copy have flipY option
-            GetTextureFormatComponentCount(dstSpec.format)};  // channelCount
+            GetTextureFormatComponentCount(dstSpec.format),  // channelCount
+            srcSpec.copyOrigin.x,
+            srcSpec.copyOrigin.y,  // src texture copy origin
+            dstSpec.copyOrigin.x,
+            dstSpec.copyOrigin.y,  // dst texture copy origin
+            copySize.width,
+            copySize.height  // copy size
+        };
 
         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.size = dstSpec.textureSize.width * dstSpec.textureSize.height * sizeof(uint32_t);
         outputDesc.usage =
             wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
         wgpu::Buffer outputBuffer = device.CreateBuffer(&outputDesc);
@@ -278,10 +360,7 @@ class CopyTextureForBrowserTests : public DawnTest {
         // Create bind group based on the config.
         wgpu::BindGroup bindGroup = utils::MakeBindGroup(
             device, testPipeline.GetBindGroupLayout(0),
-            {{0, srcTextureView},
+            {{0, srcTextureView}, {1, dstTextureView}, {2, outputBuffer}, {3, uniformBuffer}});
-             {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();
@@ -292,16 +371,18 @@ class CopyTextureForBrowserTests : public DawnTest {
             wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
             pass.SetPipeline(testPipeline);
             pass.SetBindGroup(0, bindGroup);
-            pass.Dispatch(copySize.width, copySize.height);
+            pass.Dispatch(dstSpec.textureSize.width,
+                          dstSpec.textureSize.height);  // Verify dst texture content
             pass.EndPass();
 
             testCommands = encoder.Finish();
         }
         queue.Submit(1, &testCommands);
 
-        std::vector<uint32_t> expectResult(copySize.width * copySize.height, 1);
+        std::vector<uint32_t> expectResult(dstSpec.textureSize.width * dstSpec.textureSize.height,
+                                           1);
         EXPECT_BUFFER_U32_RANGE_EQ(expectResult.data(), outputBuffer, 0,
-                                   copySize.width * copySize.height);
+                                   dstSpec.textureSize.width * dstSpec.textureSize.height);
     }
 
     wgpu::Buffer uniformBuffer;  // Uniform buffer to store dst texture meta info.
@@ -431,6 +512,35 @@ TEST_P(CopyTextureForBrowserTests, FromBGRA8UnormCopy) {
     }
 }
 
+// Verify |CopyTextureForBrowser| doing subrect copy.
+// Source texture is a full red texture and dst texture is a full
+// green texture originally. After the subrect copy, affected part
+// in dst texture should be red and other part should remain green.
+TEST_P(CopyTextureForBrowserTests, CopySubRect) {
+    // Tests skip due to crbug.com/dawn/592.
+    DAWN_SKIP_TEST_IF(IsD3D12() && IsBackendValidationEnabled());
+
+    for (wgpu::Origin3D srcOrigin : kOrigins) {
+        for (wgpu::Origin3D dstOrigin : kOrigins) {
+            for (wgpu::Extent3D copySize : kCopySize) {
+                for (bool flipY : {true, false}) {
+                    TextureSpec srcTextureSpec;
+                    srcTextureSpec.copyOrigin = srcOrigin;
+                    srcTextureSpec.textureSize = {6, 7};
+
+                    TextureSpec dstTextureSpec;
+                    dstTextureSpec.copyOrigin = dstOrigin;
+                    dstTextureSpec.textureSize = {8, 5};
+                    wgpu::CopyTextureForBrowserOptions options = {};
+                    options.flipY = flipY;
+
+                    DoTest(srcTextureSpec, dstTextureSpec, copySize, options);
+                }
+            }
+        }
+    }
+}
+
 DAWN_INSTANTIATE_TEST(CopyTextureForBrowserTests,
                       D3D12Backend(),
                       MetalBackend(),