在Metal上实现对多通道feature map的双线性插值

    kernel void resize_interpolation_array(
                                  texture2d_array<half, access::sample> inTexture [[texture(0)]],
                                  texture2d_array<half, access::write> outTexture [[texture(1)]],
                                  ushort3 gid [[thread_position_in_grid]])
    {
        if (gid.x >= outTexture.get_width() ||
            gid.y >= outTexture.get_height() ||
            gid.z >= outTexture.get_array_size()) return;

        constexpr sampler s(coord::pixel, filter::nearest, address::clamp_to_zero);

        const ushort2 pos = gid.xy;
        const ushort slice = gid.z;

        float scale_x = float(inTexture.get_width())/float(outTexture.get_width());
        float scale_y = float(inTexture.get_height())/float(outTexture.get_height());

        float i = scale_x * float(pos.x+0.5)-0.5; //this is a trick
        float j = scale_y * float(pos.y+0.5)-0.5; //this is a trick
        float i_ceil = ceil(i);
        float i_floor = floor(i);
        float j_ceil = ceil(j);
        float j_floor = floor(j);

        half4 in;
        in = inTexture.sample(s, float2(pos.x    , pos.y    ), slice);

        half4 f00 = inTexture.sample(s,float2(i_floor,j_floor),slice);
        half4 f10 = inTexture.sample(s,float2(i_ceil,j_floor),slice);
        half4 f01 = inTexture.sample(s,float2(i_floor,j_ceil),slice);
        half4 f11 = inTexture.sample(s,float2(i_ceil,j_ceil),slice);

        float4 r1 =(i_ceil-i)*float4(f01) + (i-i_floor)*float4(f11);
        float4 r2 = (i_ceil-i)*float4(f00) + (i-i_floor)*float4(f10);
        float4 out = (j-j_floor)*r1 + (j_ceil -j)*r2;

        outTexture.write(half4(out), gid.xy, gid.z);

    }

Reference:
http://handspeaker.iteye.com/blog/1545126
http://www.cnblogs.com/enigma19971/p/5828447.html
https://en.wikipedia.org/wiki/Bilinear_interpolation