dawn-cmake/test/bug/tint/534.wgsl.expected.msl
Ben Clayton 1b03f0a07a reader/wgsl: Generate ForLoopStatements
Instead of LoopStatements.

Update the writers to handle these.

Fixed: tint:952
Change-Id: Ibef66e133224810efc28c224d910b5e21f71f8d6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57203
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
2021-07-08 21:23:33 +00:00

43 lines
1.5 KiB
Plaintext

#include <metal_stdlib>
using namespace metal;
struct Uniforms {
/* 0x0000 */ uint dstTextureFlipY;
/* 0x0004 */ uint isFloat16;
/* 0x0008 */ uint isRGB10A2Unorm;
/* 0x000c */ uint channelCount;
};
struct OutputBuf {
/* 0x0000 */ uint result[1];
};
uint ConvertToFp16FloatValue(float fp32) {
return 1u;
}
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
int2 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
int2 dstTexCoord = int2(GlobalInvocationID.xy);
int2 srcTexCoord = dstTexCoord;
if ((uniforms.dstTextureFlipY == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
}
float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0);
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);
for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
success = (success && (srcColorBits[i] == dstColorBits[i]));
}
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
if (success) {
output.result[outputIndex] = uint(1);
} else {
output.result[outputIndex] = uint(0);
}
return;
}