encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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 <shaobo.yan@intel.com>
Reviewed-by: Austin Eng <enga@chromium.org>
This commit is contained in:
Yan, Shaobo 2021-03-02 03:00:27 +00:00 committed by Commit Bot service account
parent 642ec1c30e
commit fbda46da44
1 changed files with 141 additions and 87 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

228
src/tests/end2end/CopyTextureForBrowserTests.cpp
View File

@ -16,6 +16,7 @@
#include "common/Constants.h" #include "common/Constants.h"
#include "common/Math.h" #include "common/Math.h"
#include "utils/ComboRenderPipelineDescriptor.h"
#include "utils/TestUtils.h" #include "utils/TestUtils.h"
#include "utils/TextureFormatUtils.h" #include "utils/TextureFormatUtils.h"
#include "utils/WGPUHelpers.h" #include "utils/WGPUHelpers.h"
@ -23,14 +24,17 @@
class CopyTextureForBrowserTests : public DawnTest { class CopyTextureForBrowserTests : public DawnTest {
protected: protected:
static constexpr wgpu::TextureFormat kTextureFormat = wgpu::TextureFormat::RGBA8Unorm; static constexpr wgpu::TextureFormat kTextureFormat = wgpu::TextureFormat::RGBA8Unorm;
static constexpr uint64_t kDefaultTextureWidth = 4;
static constexpr uint64_t kDefaultTextureHeight = 4;
struct TextureSpec { struct TextureSpec {
wgpu::Origin3D copyOrigin; wgpu::Origin3D copyOrigin = {};
wgpu::Extent3D textureSize; wgpu::Extent3D textureSize = {kDefaultTextureWidth, kDefaultTextureHeight};
uint32_t level; uint32_t level = 0;
wgpu::TextureFormat format = kTextureFormat;
}; };
static std::vector<RGBA8> GetExpectedTextureData(const utils::TextureDataCopyLayout& layout) { static std::vector<RGBA8> GetSourceTextureData(const utils::TextureDataCopyLayout& layout) {
std::vector<RGBA8> textureData(layout.texelBlockCount); std::vector<RGBA8> textureData(layout.texelBlockCount);
for (uint32_t layer = 0; layer < layout.mipSize.depth; ++layer) { for (uint32_t layer = 0; layer < layout.mipSize.depth; ++layer) {
const uint32_t sliceOffset = layout.texelBlocksPerImage * layer; const uint32_t sliceOffset = layout.texelBlocksPerImage * layer;
@ -48,43 +52,87 @@ class CopyTextureForBrowserTests : public DawnTest {
return textureData; return textureData;
} }
static void PackTextureData(const RGBA8* srcData, void SetUp() override {
uint32_t width, DawnTest::SetUp();
uint32_t height,
uint32_t srcTexelsPerRow, testPipeline = MakeTestPipeline();
RGBA8* dstData,
uint32_t dstTexelsPerRow, uint32_t uniformBufferData[] = {
const wgpu::CopyTextureForBrowserOptions* options) { 0, // copy have flipY option
bool isFlipY = options != nullptr && options->flipY; };
for (uint32_t y = 0; y < height; ++y) {
for (uint32_t x = 0; x < width; ++x) { wgpu::BufferDescriptor uniformBufferDesc = {};
uint32_t srcYIndex = uniformBufferDesc.usage = wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Uniform;
isFlipY ? (height - y - 1) * srcTexelsPerRow : y * srcTexelsPerRow; uniformBufferDesc.size = sizeof(uniformBufferData);
uint32_t src = x + srcYIndex; uniformBuffer = device.CreateBuffer(&uniformBufferDesc);
uint32_t dst = x + y * dstTexelsPerRow; }
dstData[dst] = srcData[src];
// 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<u32>;
};
[[group(0), binding(0)]] var src : texture_2d<f32>;
[[group(0), binding(1)]] var dst : texture_2d<f32>;
[[group(0), binding(2)]] var<storage_buffer> output : OutputBuf;
[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
[[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<i32> = textureDimensions(src);
var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
var srcTexCoord : vec2<i32> = dstTexCoord;
if (uniforms.dstTextureFlipY == 1) {
srcTexCoord.y = size.y - dstTexCoord.y - 1;
}
var srcColor : vec4<f32> = textureLoad(src, srcTexCoord);
var dstColor : vec4<f32> = 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, void DoTest(const TextureSpec& srcSpec,
const TextureSpec& dstSpec, const TextureSpec& dstSpec,
const wgpu::Extent3D& copySize, const wgpu::Extent3D& copySize = {kDefaultTextureWidth, kDefaultTextureHeight},
const wgpu::CopyTextureForBrowserOptions* options) { const wgpu::CopyTextureForBrowserOptions options = {}) {
wgpu::TextureDescriptor srcDescriptor; wgpu::TextureDescriptor srcDescriptor;
srcDescriptor.size = srcSpec.textureSize; srcDescriptor.size = srcSpec.textureSize;
srcDescriptor.format = kTextureFormat; srcDescriptor.format = srcSpec.format;
srcDescriptor.mipLevelCount = srcSpec.level + 1; srcDescriptor.mipLevelCount = srcSpec.level + 1;
srcDescriptor.usage = wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | srcDescriptor.usage =
wgpu::TextureUsage::Sampled | wgpu::TextureUsage::OutputAttachment; wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | wgpu::TextureUsage::Sampled;
wgpu::Texture srcTexture = device.CreateTexture(&srcDescriptor); wgpu::Texture srcTexture = device.CreateTexture(&srcDescriptor);
wgpu::Texture dstTexture; wgpu::Texture dstTexture;
wgpu::TextureDescriptor dstDescriptor; wgpu::TextureDescriptor dstDescriptor;
dstDescriptor.size = dstSpec.textureSize; dstDescriptor.size = dstSpec.textureSize;
dstDescriptor.format = kTextureFormat; dstDescriptor.format = dstSpec.format;
dstDescriptor.mipLevelCount = dstSpec.level + 1; dstDescriptor.mipLevelCount = dstSpec.level + 1;
dstDescriptor.usage = wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst | dstDescriptor.usage = wgpu::TextureUsage::CopyDst | wgpu::TextureUsage::Sampled |
wgpu::TextureUsage::OutputAttachment; wgpu::TextureUsage::OutputAttachment;
dstTexture = device.CreateTexture(&dstDescriptor); dstTexture = device.CreateTexture(&dstDescriptor);
@ -96,7 +144,7 @@ class CopyTextureForBrowserTests : public DawnTest {
{srcSpec.textureSize.width, srcSpec.textureSize.height, copySize.depth}, {srcSpec.textureSize.width, srcSpec.textureSize.height, copySize.depth},
srcSpec.level); srcSpec.level);
const std::vector<RGBA8> textureArrayCopyData = GetExpectedTextureData(copyLayout); const std::vector<RGBA8> textureArrayCopyData = GetSourceTextureData(copyLayout);
wgpu::TextureCopyView textureCopyView = wgpu::TextureCopyView textureCopyView =
utils::CreateTextureCopyView(srcTexture, srcSpec.level, {0, 0, srcSpec.copyOrigin.z}); utils::CreateTextureCopyView(srcTexture, srcSpec.level, {0, 0, srcSpec.copyOrigin.z});
@ -109,7 +157,6 @@ class CopyTextureForBrowserTests : public DawnTest {
textureArrayCopyData.size() * sizeof(RGBA8), textureArrayCopyData.size() * sizeof(RGBA8),
&textureDataLayout, &copyLayout.mipSize); &textureDataLayout, &copyLayout.mipSize);
const wgpu::Extent3D copySizePerSlice = {copySize.width, copySize.height, 1};
// Perform the texture to texture copy // Perform the texture to texture copy
wgpu::TextureCopyView srcTextureCopyView = wgpu::TextureCopyView srcTextureCopyView =
utils::CreateTextureCopyView(srcTexture, srcSpec.level, srcSpec.copyOrigin); utils::CreateTextureCopyView(srcTexture, srcSpec.level, srcSpec.copyOrigin);
@ -119,41 +166,64 @@ class CopyTextureForBrowserTests : public DawnTest {
wgpu::CommandBuffer commands = encoder.Finish(); wgpu::CommandBuffer commands = encoder.Finish();
queue.Submit(1, &commands); queue.Submit(1, &commands);
// Perform a copy here for testing.
device.GetQueue().CopyTextureForBrowser(&srcTextureCopyView, &dstTextureCopyView, &copySize, device.GetQueue().CopyTextureForBrowser(&srcTextureCopyView, &dstTextureCopyView, &copySize,
options); &options);
// Texels in single slice. // Update uniform buffer based on test config
const uint32_t texelCountInCopyRegion = utils::GetTexelCountInCopyRegion( uint32_t uniformBufferData[] = {
copyLayout.bytesPerRow, copyLayout.bytesPerImage / copyLayout.bytesPerRow, options.flipY, // copy have flipY option
copySizePerSlice, kTextureFormat); };
std::vector<RGBA8> 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);
EXPECT_TEXTURE_RGBA8_EQ(expected.data(), dstTexture, dstSpec.copyOrigin.x, device.GetQueue().WriteBuffer(uniformBuffer, 0, uniformBufferData,
dstSpec.copyOrigin.y, copySize.width, copySize.height, sizeof(uniformBufferData));
dstSpec.level, dstSpec.copyOrigin.z + slice)
<< "Texture to Texture copy failed copying region [(" << srcSpec.copyOrigin.x // Create output buffer to store result
<< ", " << srcSpec.copyOrigin.y << "), (" << srcSpec.copyOrigin.x + copySize.width wgpu::BufferDescriptor outputDesc;
<< ", " << srcSpec.copyOrigin.y + copySize.height << ")) from " outputDesc.size = copySize.width * copySize.height * sizeof(uint32_t);
<< srcSpec.textureSize.width << " x " << srcSpec.textureSize.height outputDesc.usage =
<< " texture at mip level " << srcSpec.level << " layer " wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
<< srcSpec.copyOrigin.z + slice << " to [(" << dstSpec.copyOrigin.x << ", " wgpu::Buffer outputBuffer = device.CreateBuffer(&outputDesc);
<< dstSpec.copyOrigin.y << "), (" << dstSpec.copyOrigin.x + copySize.width << ", "
<< dstSpec.copyOrigin.y + copySize.height << ")) region of " // Create texture views for test.
<< dstSpec.textureSize.width << " x " << dstSpec.textureSize.height wgpu::TextureViewDescriptor srcTextureViewDesc = {};
<< " texture at mip level " << dstSpec.level << " layer " srcTextureViewDesc.baseMipLevel = srcSpec.level;
<< dstSpec.copyOrigin.z + slice << std::endl; 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<uint32_t> 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. // Verify CopyTextureForBrowserTests works with internal pipeline.
@ -166,12 +236,9 @@ TEST_P(CopyTextureForBrowserTests, PassthroughCopy) {
constexpr uint32_t kHeight = 1; constexpr uint32_t kHeight = 1;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; DoTest(textureSpec, textureSpec, {kWidth, kHeight});
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options);
} }
TEST_P(CopyTextureForBrowserTests, VerifyCopyOnXDirection) { TEST_P(CopyTextureForBrowserTests, VerifyCopyOnXDirection) {
@ -182,12 +249,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyOnXDirection) {
constexpr uint32_t kHeight = 1; constexpr uint32_t kHeight = 1;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; DoTest(textureSpec, textureSpec, {kWidth, kHeight});
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options);
} }
TEST_P(CopyTextureForBrowserTests, VerifyCopyOnYDirection) { TEST_P(CopyTextureForBrowserTests, VerifyCopyOnYDirection) {
@ -198,12 +262,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyOnYDirection) {
constexpr uint32_t kHeight = 1000; constexpr uint32_t kHeight = 1000;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; DoTest(textureSpec, textureSpec, {kWidth, kHeight});
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options);
} }
TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) { TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) {
@ -214,12 +275,9 @@ TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) {
constexpr uint32_t kHeight = 999; constexpr uint32_t kHeight = 999;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; DoTest(textureSpec, textureSpec, {kWidth, kHeight});
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options);
} }
TEST_P(CopyTextureForBrowserTests, VerifyFlipY) { TEST_P(CopyTextureForBrowserTests, VerifyFlipY) {
@ -230,13 +288,11 @@ TEST_P(CopyTextureForBrowserTests, VerifyFlipY) {
constexpr uint32_t kHeight = 1001; constexpr uint32_t kHeight = 1001;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; wgpu::CopyTextureForBrowserOptions options = {};
options.flipY = true; options.flipY = true;
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); DoTest(textureSpec, textureSpec, {kWidth, kHeight}, options);
} }
TEST_P(CopyTextureForBrowserTests, VerifyFlipYInSlimTexture) { TEST_P(CopyTextureForBrowserTests, VerifyFlipYInSlimTexture) {
@ -247,13 +303,11 @@ TEST_P(CopyTextureForBrowserTests, VerifyFlipYInSlimTexture) {
constexpr uint32_t kHeight = 1001; constexpr uint32_t kHeight = 1001;
TextureSpec textureSpec; TextureSpec textureSpec;
textureSpec.copyOrigin = {0, 0, 0}; textureSpec.textureSize = {kWidth, kHeight};
textureSpec.level = 0;
textureSpec.textureSize = {kWidth, kHeight, 1};
wgpu::CopyTextureForBrowserOptions options = {}; wgpu::CopyTextureForBrowserOptions options = {};
options.flipY = true; options.flipY = true;
DoTest(textureSpec, textureSpec, {kWidth, kHeight, 1}, &options); DoTest(textureSpec, textureSpec, {kWidth, kHeight}, options);
} }
DAWN_INSTANTIATE_TEST(CopyTextureForBrowserTests, DAWN_INSTANTIATE_TEST(CopyTextureForBrowserTests,