CUDA bilinear resize

__global__ void bilinear_resize(unsigned int *out, int width, float ratio) {
    int x, y;
    int j = blockDim.x * blockIdx.x + threadIdx.x;
    int i = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int a, b, c, d;
    float x_diff, y_diff, blue, red, green;

    // do the bilinear scaling
    x = (int)(ratio * j);
    y = (int)(ratio * i);
    x_diff = (ratio * j) - x;
    y_diff = (ratio * i) - y;

    // fetch the 32bit ARGB unsigned ints for the four pixels
    a = tex2D(texRef, x, y);
    b = tex2D(texRef, x+1, y);
    c = tex2D(texRef, x, y+1);
    d = tex2D(texRef, x+1, y+1);

    // blue channel
    // Yb = Ab(1-w)(1-h) + Bb(w)(1-h) + Cb(h)(1-w) + Db(wh)
    blue = (MAKE_BLUE_PIXEL(a) * (1.0f-x_diff) * (1.0f-y_diff)) +
           (MAKE_BLUE_PIXEL(b) * (x_diff) * (1.0f-y_diff)) +
           (MAKE_BLUE_PIXEL(c) * (1.0f-x_diff) * (y_diff)) +
           (MAKE_BLUE_PIXEL(d) * (x_diff * y_diff));

    // green channel
    // Yg = Ag(1-w)(1-h) + Bg(w)(1-h) + Cg(h)(1-w) + Dg(wh)
    green = (MAKE_GREEN_PIXEL(a) * (1.0f-x_diff) * (1.0f-y_diff)) +
            (MAKE_GREEN_PIXEL(b) * (x_diff) * (1.0f-y_diff)) +
            (MAKE_GREEN_PIXEL(c) * (1.0f-x_diff) * (y_diff)) +
            (MAKE_GREEN_PIXEL(d) * (x_diff * y_diff));

    // red channel
    // Yr = Ar(1-w)(1-h) + Br(w)(1-h) + Cr(h)(1-w) + Dr(wh)
    red = (MAKE_RED_PIXEL(a) * (1.0f-x_diff) * (1.0f-y_diff)) +
          (MAKE_RED_PIXEL(b) * (x_diff) * (1.0f-y_diff)) +
          (MAKE_RED_PIXEL(c) * (1.0f-x_diff) * (y_diff)) +
          (MAKE_RED_PIXEL(d) * (x_diff * y_diff));

    out[i*width + j] = MAKE_ARGB(
        CLAMP(__float2uint_rn(red), 0, 255),
        CLAMP(__float2uint_rn(green), 0, 255),
        CLAMP(__float2uint_rn(blue), 0, 255));
}