[[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;
    }
}