diff --git a/src/dawn_native/CopyTextureForBrowserHelper.cpp b/src/dawn_native/CopyTextureForBrowserHelper.cpp index b6374bec2b..1dc232989d 100644 --- a/src/dawn_native/CopyTextureForBrowserHelper.cpp +++ b/src/dawn_native/CopyTextureForBrowserHelper.cpp @@ -49,10 +49,25 @@ namespace dawn_native { [[stage(vertex)]] fn main() -> void { Position = vec4((texcoord[VertexIndex] * 2.0 - vec2(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(1.0, -1.0) + vec2(0.0, 1.0)) * - uniforms.u_scale + uniforms.u_offset; + if (flipY) { + // 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(1.0, -1.0) + vec2(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(1.0, -1.0) + vec2(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 - 0.0, 0.0 // offset + copySize->width / static_cast(srcTextureSize.width), + copySize->height / static_cast(srcTextureSize.height), // scale + source->origin.x / static_cast(srcTextureSize.width), + source->origin.y / static_cast(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(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(); diff --git a/src/tests/end2end/CopyTextureForBrowserTests.cpp b/src/tests/end2end/CopyTextureForBrowserTests.cpp index 33df4a0441..8f838fdfea 100644 --- a/src/tests/end2end/CopyTextureForBrowserTests.cpp +++ b/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 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 GetTextureData(const utils::TextureDataCopyLayout& layout, + TextureCopyRole textureRole) { std::vector 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] = - RGBA8(static_cast((x + layer * x) % 256), - static_cast((y + layer * y) % 256), - static_cast(x % 256), static_cast(x % 256)); + // Source textures will have variable pixel data to cover cases like + // flipY. + if (textureRole == TextureCopyRole::SOURCE) { + textureData[sliceOffset + rowOffset + x] = + RGBA8(static_cast((x + layer * x) % 256), + static_cast((y + layer * y) % 256), + static_cast(x % 256), static_cast(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(0), static_cast(255), + static_cast(0), static_cast(255)); + } } } } @@ -99,8 +119,11 @@ class CopyTextureForBrowserTests : public DawnTest { testPipeline = MakeTestPipeline(); uint32_t uniformBufferData[] = { - 0, // copy have flipY option - 4, // channelCount + 0, // copy have flipY option + 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; + dstCopyOrigin : vec2; + copySize : vec2; }; [[block]] struct OutputBuf { result : array; @@ -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. - // TODO(crbug.com/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 == 1u) { - srcTexCoord.y = size.y - dstTexCoord.y - 1; - } + var srcSize : vec2 = textureDimensions(src); + var dstSize : vec2 = textureDimensions(dst); + var dstTexCoord : vec2 = vec2(GlobalInvocationID.xy); + var nonCoveredColor : vec4 = + vec4(0.0, 1.0, 0.0, 1.0); // should be green - var srcColor : vec4 = textureLoad(src, srcTexCoord, 0); - var dstColor : vec4 = textureLoad(dst, dstTexCoord, 0); var success : bool = true; - - // Not use loop and variable index format to workaround - // crbug.com/tint/638. - if (uniforms.channelCount == 2u) { // All have rg components. + if (dstTexCoord.x < uniforms.dstCopyOrigin.x || + dstTexCoord.y < uniforms.dstCopyOrigin.y || + dstTexCoord.x >= uniforms.dstCopyOrigin.x + uniforms.copySize.x || + dstTexCoord.y >= uniforms.dstCopyOrigin.y + uniforms.copySize.y) { success = success && - aboutEqual(dstColor.r, srcColor.r) && - aboutEqual(dstColor.g, srcColor.g); + all(textureLoad(dst, vec2(dstTexCoord), 0) == nonCoveredColor); } else { - success = success && - aboutEqual(dstColor.r, srcColor.r) && - aboutEqual(dstColor.g, srcColor.g) && - aboutEqual(dstColor.b, srcColor.b) && - aboutEqual(dstColor.a, srcColor.a); - } + // Calculate source texture coord. + var srcTexCoord : vec2 = dstTexCoord - uniforms.dstCopyOrigin + + uniforms.srcCopyOrigin; + // Note that |flipY| equals flip src texture firstly and then do copy from src + // 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 = textureLoad(src, vec2(srcTexCoord), 0); + var dstColor : vec4 = textureLoad(dst, vec2(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 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 = - utils::GetTextureDataCopyLayoutForTexture2DAtLevel( - kTextureFormat, - {srcSpec.textureSize.width, srcSpec.textureSize.height, - copySize.depthOrArrayLayers}, - srcSpec.level); + const std::vector dstTextureArrayCopyData = + GetTextureData(dstCopyLayout, TextureCopyRole::DEST); - const std::vector textureArrayCopyData = - useFixedTestValue ? GetFixedSourceTextureData() : GetSourceTextureData(copyLayout); - wgpu::ImageCopyTexture imageCopyTexture = - utils::CreateImageCopyTexture(srcTexture, srcSpec.level, {0, 0, srcSpec.copyOrigin.z}); + wgpu::TextureDataLayout dstTextureDataLayout; + dstTextureDataLayout.offset = 0; + dstTextureDataLayout.bytesPerRow = dstCopyLayout.bytesPerRow; + dstTextureDataLayout.rowsPerImage = dstCopyLayout.rowsPerImage; - wgpu::TextureDataLayout textureDataLayout; - textureDataLayout.offset = 0; - textureDataLayout.bytesPerRow = copyLayout.bytesPerRow; - textureDataLayout.rowsPerImage = copyLayout.rowsPerImage; + wgpu::ImageCopyTexture dstImageTextureInit = + utils::CreateImageCopyTexture(dstTexture, dstSpec.level, {0, 0}); - device.GetQueue().WriteTexture(&imageCopyTexture, textureArrayCopyData.data(), - textureArrayCopyData.size() * sizeof(RGBA8), - &textureDataLayout, ©Layout.mipSize); + device.GetQueue().WriteTexture(&dstImageTextureInit, dstTextureArrayCopyData.data(), + dstTextureArrayCopyData.size() * sizeof(RGBA8), + &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, ©Size, &options); // Update uniform buffer based on test config uint32_t uniformBufferData[] = { - options.flipY, // copy have flipY option - GetTextureFormatComponentCount(dstSpec.format)}; // channelCount + options.flipY, // copy have flipY option + 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}, - {1, dstTextureView}, - {2, outputBuffer, 0, copySize.width * copySize.height * sizeof(uint32_t)}, - {3, uniformBuffer, 0, sizeof(uniformBufferData)}}); + {{0, srcTextureView}, {1, dstTextureView}, {2, outputBuffer}, {3, uniformBuffer}}); // 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 expectResult(copySize.width * copySize.height, 1); + std::vector 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(),