#include <metal_stdlib>
using namespace metal;

struct BlitParams
{
    uint3 offset;
    uint3 size;
    uint2 stride;
};

kernel void blitTexture3D_RGBA8(
  uint3 gid [[ thread_position_in_grid ]],
  constant BlitParams& params [[ buffer(0) ]],
  device uchar4* input [[ buffer(1) ]],
  texture3d<float, access::write> output [[ texture(0) ]]
  )
{
    if (all(gid < params.size))
    {
        uchar4 element = input[gid.x + gid.y * params.stride.x + gid.z * params.stride.y];

        output.write(float4(element) / 255.f, gid + params.offset);
    }
}

kernel void blitTexture2D_R8(
  uint3 gid [[ thread_position_in_grid ]],
  constant BlitParams& params [[ buffer(0) ]],
  device uchar* input [[ buffer(1) ]],
  texture2d<float, access::write> output [[ texture(0) ]]
  )
{
    if (all(gid < params.size))
    {
        uchar element = input[gid.x + gid.y * params.stride.x];

        output.write(float(element) / 255.f, gid.xy + params.offset.xy);
    }
}

kernel void blitTexture3D_RG8(
  uint3 gid [[ thread_position_in_grid ]],
  constant BlitParams& params [[ buffer(0) ]],
  device uchar2* input [[ buffer(1) ]],
  texture3d<float, access::write> output [[ texture(0) ]]
  )
{
    if (all(gid < params.size))
    {
        uchar2 element = input[gid.x + gid.y * params.stride.x + gid.z * params.stride.y];

        output.write(float4(float2(element) / 255.f, 0.f, 0.f), gid + params.offset);
    }
}
