blob: 467a9443a040a0337d459aee138cb13878176279 [file] [log] [blame]
[[block]] struct Uniforms {
dstTextureFlipY : u32;
channelCount : u32;
srcCopyOrigin : vec2<u32>;
dstCopyOrigin : vec2<u32>;
copySize : vec2<u32>;
};
[[block]] struct OutputBuf {
result : 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, read_write> output : OutputBuf;
[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
fn aboutEqual(value : f32, expect : f32) -> bool {
// The value diff should be smaller than the hard coded tolerance.
return abs(value - expect) < 0.001;
}
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
let srcSize : vec2<i32> = textureDimensions(src);
let dstSize : vec2<i32> = textureDimensions(dst);
let dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
let nonCoveredColor : vec4<f32> =
vec4<f32>(0.0, 1.0, 0.0, 1.0); // should be green
var success : bool = true;
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 &&
all(textureLoad(dst, vec2<i32>(dstTexCoord), 0) == nonCoveredColor);
} else {
// 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;
}
let srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
let 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);
}
}
let outputIndex : u32 = GlobalInvocationID.y * u32(dstSize.x) +
GlobalInvocationID.x;
if (success) {
output.result[outputIndex] = 1u;
} else {
output.result[outputIndex] = 0u;
}
}