CL PNG filter

//Image border color must be 0x00000000
//Filtered result is 32bpp RGBA.
__kernel void png_filterRow(__global __read_only image2d_t raw,
                            __global uchar *filtered,
                            __local uint4 *rcache,
                            __local uint4 *ucache) {
    //Image sampler for reading raw
    const sampler_t ss = CLK_NORMALIZED_COORDS_FALSE |  //Use pixel coords
                         CLK_ADDRESS_CLAMP |            //Clamp to black
                         CLK_FILTER_NEAREST;            //No interpolation

    //One work-group per image row, one worker per pixel
    size_t row = get_group_id(0);
    size_t rowlen = get_image_width(raw);
    size_t col = get_local_id(0);

    //Do something to figure which filter type to use for the row
    uchar filter_type = PNG_FILTER_PAETH;
    //Assuming 24-bit color, keeping the int4 to avoid saturating filters
    uint4 pixel = read_imageui(raw, ss, (int2)(col, row));
    //TODO: filtered values are probably signed.  Verify that & fix

    //Cache adjacent pixels in __local if useful for the filter
    if (filter_type != PNG_FILTER_NONE) {
        //Needs adjusting if > 1 workgroup per row
        //TODO: check image addressing- is (0,0) upper or lower left?
        rcache[0] = read_imageui(raw, ss, (int2)(-1,row));
        rcache[col+1] = pixel;
        if (filter_type != PNG_FILTER_SUB) {
            //Up, Average and Paeth all look at preceding row too, cache that
            ucache[0] = read_imageui(raw, ss, (int2)(-1, row-1));
            ucache[col+1] = read_imageui(raw, ss, (int2)(col, row-1)
        }
        barrier(CLK_LOCAL_MEM_FENCE);       //Sync {u,r}cache
    }

    uint4 predicted;
    switch (filter_type) {
        case PNG_FILTER_NONE:
            predicted = 0;
            break;
        case PNG_FILTER_SUB:
            predicted = rcache[col];
            break;
        case PNG_FILTER_UP:
            predicted = ucache[col+1];
            break;
        case PNG_FILTER_AVERAGE:
            predicted = (rcache[col] + ucache[col+1]) / 2;
            break;
        case PNG_FILTER_PAETH:
            uint4 a = rcache[col], b = ucache[col+1], c = ucache[col];
            uint4 p = a + b - c;
            //Hopefully minmag will shuffle members to select the nearest
            //component from each vector, else need nasty logic.
            predicted = minmag(a-p, b-p);
            predicted = minmag(predicted, c-p);
            break;
    }

    //Write out filtered pixel
    pixel -= predicted;
    uchar4 bytes = convert_uchar4(pixel);
    vstore4(bytes, (row * rowlen) + col, filtered);
}