[[block]] struct Uniforms {
    dstTextureFlipY : u32;
    isFloat16 : u32;
    isRGB10A2Unorm : u32;
    channelCount : u32;
};
[[block]] struct OutputBuf {
    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, read_write> output : OutputBuf;
[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
//[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
// Fp16 logic
// Infinity and NaN won't happen in this test case.
fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
    return 1u;
}

[[stage(compute), workgroup_size(1, 1, 1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
    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 srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
    var dstColor : vec4<f32> = textureLoad(dst, dstTexCoord, 0);
    var success : bool = true;

    var srcColorBits : vec4<u32>;
    var dstColorBits : vec4<u32> = vec4<u32>(dstColor);

    for (var i : u32 = 0u; i < uniforms.channelCount; i = i + 1u) {
            srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
            success = success && (srcColorBits[i] == dstColorBits[i]);
    }

    var outputIndex : u32 = GlobalInvocationID.y * u32(size.x) + GlobalInvocationID.x;
    if (success) {
        output.result[outputIndex] = u32(1);
    } else {
        output.result[outputIndex] = u32(0);
    }
}