#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

kernel void main0(texture2d<float> uImageIn [[texture(0)]], texture2d<float, access::write> uImageOut [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    int2 _23 = int2(gl_GlobalInvocationID.xy);
    uImageOut.write(uImageIn.read(uint2((_23 + int2(uImageIn.get_width(), uImageIn.get_height())))), uint2(_23));
}