test/tint/bug/tint/534.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;

template<typename T, size_t N>
struct tint_array {
    const constant T& operator[](size_t i) const constant { return elements[i]; }
    device T& operator[](size_t i) device { return elements[i]; }
    const device T& operator[](size_t i) const device { return elements[i]; }
    thread T& operator[](size_t i) thread { return elements[i]; }
    const thread T& operator[](size_t i) const thread { return elements[i]; }
    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
    T elements[N];
};

#define TINT_ISOLATE_UB(VOLATILE_NAME) \
  volatile bool VOLATILE_NAME = true; \
  if (VOLATILE_NAME)

uint4 tint_ftou(float4 v) {
  return select(uint4(4294967295u), select(uint4(v), uint4(0u), (v < float4(0.0f))), (v <= float4(4294967040.0f)));
}

struct Uniforms {
  /* 0x0000 */ uint dstTextureFlipY;
  /* 0x0004 */ uint isFloat16;
  /* 0x0008 */ uint isRGB10A2Unorm;
  /* 0x000c */ uint channelCount;
};

struct OutputBuf {
  /* 0x0000 */ tint_array<uint, 1> result;
};

uint ConvertToFp16FloatValue(float fp32) {
  return 1u;
}

void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_2, const constant Uniforms* const tint_symbol_3, texture2d<float, access::sample> tint_symbol_4, device OutputBuf* const tint_symbol_5) {
  uint2 size = uint2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
  uint2 dstTexCoord = GlobalInvocationID.xy;
  uint2 srcTexCoord = dstTexCoord;
  if (((*(tint_symbol_3)).dstTextureFlipY == 1u)) {
    srcTexCoord[1] = ((size[1] - dstTexCoord[1]) - 1u);
  }
  float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
  float4 dstColor = tint_symbol_4.read(uint2(dstTexCoord), 0);
  bool success = true;
  uint4 srcColorBits = 0u;
  uint4 dstColorBits = tint_ftou(dstColor);
  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
    uint const tint_symbol_1 = i;
    srcColorBits[tint_symbol_1] = ConvertToFp16FloatValue(srcColor[i]);
    success = (success && (srcColorBits[i] == dstColorBits[i]));
  }
  uint outputIndex = ((GlobalInvocationID[1] * uint(size[0])) + GlobalInvocationID[0]);
  if (success) {
    (*(tint_symbol_5)).result[outputIndex] = 1u;
  } else {
    (*(tint_symbol_5)).result[outputIndex] = 0u;
  }
}

kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_6 [[texture(0)]], const constant Uniforms* tint_symbol_7 [[buffer(0)]], texture2d<float, access::sample> tint_symbol_8 [[texture(1)]], device OutputBuf* tint_symbol_9 [[buffer(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
  tint_symbol_inner(GlobalInvocationID, tint_symbol_6, tint_symbol_7, tint_symbol_8, tint_symbol_9);
  return;
}