2021-06-30 15:06:00 +00:00
|
|
|
struct Uniforms {
|
|
|
|
dstTextureFlipY : u32;
|
|
|
|
isFloat16 : u32;
|
|
|
|
isRGB10A2Unorm : u32;
|
|
|
|
channelCount : u32;
|
2022-01-19 18:11:17 +00:00
|
|
|
}
|
2021-06-30 15:06:00 +00:00
|
|
|
|
|
|
|
struct OutputBuf {
|
2022-01-20 22:11:07 +00:00
|
|
|
result : array<u32>;
|
2022-01-19 18:11:17 +00:00
|
|
|
}
|
2021-06-30 15:06:00 +00:00
|
|
|
|
2022-01-19 22:46:57 +00:00
|
|
|
@group(0) @binding(0) var src : texture_2d<f32>;
|
2021-06-30 15:06:00 +00:00
|
|
|
|
2022-01-19 22:46:57 +00:00
|
|
|
@group(0) @binding(1) var dst : texture_2d<f32>;
|
2021-06-30 15:06:00 +00:00
|
|
|
|
2022-01-19 22:46:57 +00:00
|
|
|
@group(0) @binding(2) var<storage, read_write> output : OutputBuf;
|
2021-06-30 15:06:00 +00:00
|
|
|
|
2022-01-19 22:46:57 +00:00
|
|
|
@group(0) @binding(3) var<uniform> uniforms : Uniforms;
|
2021-06-30 15:06:00 +00:00
|
|
|
|
|
|
|
fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
|
|
|
|
return 1u;
|
|
|
|
}
|
|
|
|
|
2022-01-19 22:46:57 +00:00
|
|
|
@stage(compute) @workgroup_size(1, 1, 1)
|
|
|
|
fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
|
2021-06-30 15:06:00 +00:00
|
|
|
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);
|
2021-07-08 21:23:33 +00:00
|
|
|
for(var i : u32 = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
|
|
|
|
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
|
|
|
|
success = (success && (srcColorBits[i] == dstColorBits[i]));
|
2021-06-30 15:06:00 +00:00
|
|
|
}
|
|
|
|
var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
|
|
|
|
if (success) {
|
|
|
|
output.result[outputIndex] = u32(1);
|
|
|
|
} else {
|
|
|
|
output.result[outputIndex] = u32(0);
|
|
|
|
}
|
|
|
|
}
|