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>
diff --git a/src/dawn_native/CopyTextureForBrowserHelper.cpp b/src/dawn_native/CopyTextureForBrowserHelper.cpp
index b6374be..1dc2329 100644
--- a/src/dawn_native/CopyTextureForBrowserHelper.cpp
+++ b/src/dawn_native/CopyTextureForBrowserHelper.cpp
@@ -49,10 +49,25 @@
[[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)) *
- 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<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 @@
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 @@
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<float>(srcTextureSize.width),
+ 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 @@
// 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 33df4a0..8f838fd 100644
--- a/src/tests/end2end/CopyTextureForBrowserTests.cpp
+++ b/src/tests/end2end/CopyTextureForBrowserTests.cpp
@@ -34,6 +34,10 @@
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 @@
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] =
- 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));
+ // Source textures will have variable pixel data to cover cases like
+ // flipY.
+ if (textureRole == TextureCopyRole::SOURCE) {
+ 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 @@
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 @@
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 @@
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<i32> = textureDimensions(src);
- var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
- var srcTexCoord : vec2<i32> = dstTexCoord;
- if (uniforms.dstTextureFlipY == 1u) {
- srcTexCoord.y = size.y - dstTexCoord.y - 1;
- }
+ var srcSize : vec2<i32> = textureDimensions(src);
+ var dstSize : vec2<i32> = textureDimensions(dst);
+ var dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
+ var nonCoveredColor : vec4<f32> =
+ vec4<f32>(0.0, 1.0, 0.0, 1.0); // should be green
- var srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
- var dstColor : vec4<f32> = 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<i32>(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<u32> = 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<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 @@
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 @@
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 @@
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<RGBA8> dstTextureArrayCopyData =
+ GetTextureData(dstCopyLayout, TextureCopyRole::DEST);
- const std::vector<RGBA8> 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 @@
// 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 @@
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 @@
}
}
+// 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(),
