cl_kernel/kernel_wavelet_denoise.cl

/*
 * function: kernel_wavelet_denoise
 *     wavelet filter for denoise usage
 * in:        input image data as read only
 * threshold:   noise threshold
 * low:
 */

__constant float threshConst[5] = { 50.430166f, 20.376415f, 10.184031f, 6.640919f, 3.367972f };

__kernel void kernel_wavelet_denoise(__global uint *src, __global uint *approxOut, __global float *details, __global uint *dest,
                                     int inputYOffset, int outputYOffset, uint inputUVOffset, uint outputUVOffset,
                                     int layer, int decomLevels, float hardThresh, float softThresh)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    size_t width = get_global_size(0);
    size_t height = get_global_size(1);

    int imageWidth = width * 16;
    int imageHeight = height;

    float stdev = 0.0f;
    float thold = 0.0f;
    float16 deviation = (float16)0.0f;

    layer = (layer > 1) ? layer : 1;
    layer = (layer < decomLevels) ? layer : decomLevels;

    src += inputYOffset;
    dest += outputYOffset;

#if WAVELET_DENOISE_UV
    int xScaler = pown(2.0f, layer);
    int yScaler = pown(2.0f, (layer - 1));
#else
    int xScaler = pown(2.0f, (layer - 1));
    int yScaler = xScaler;
#endif

    xScaler = ((x == 0) || (x > imageWidth / 16 - xScaler)) ? 0 : xScaler;
    yScaler = ((y < yScaler) || (y > imageHeight - yScaler)) ? 0 : yScaler;

    uint4 approx;
    float16 detail;

#if WAVELET_DENOISE_UV
    int srcOffset = (layer % 2) ? (inputUVOffset * imageWidth / 4) : 0;
    __global uchar *src_p = (__global uchar *)(src + srcOffset);
#else
    __global uchar *src_p = (__global uchar *)(src);
#endif

    int pixel_index = x * 16 + y * imageWidth;
    int group_index = x * 4 + y * (imageWidth / 4);

#if WAVELET_DENOISE_UV
    uint4 luma;
    int luma_index0 = x * 4 + (2 * y) * (imageWidth / 4);
    int luma_index1 = x * 4 + (2 * y + 1) * (imageWidth / 4);
#else
    uint4 chroma;
    int chroma_index = x * 4 + (y / 2) * (imageWidth / 4);
#endif

    ushort16 a;
    ushort16 b;
    ushort16 c;
    ushort16 d;
    ushort16 e;
    ushort16 f;
    ushort16 g;
    ushort16 h;
    ushort16 i;

    float div = 1.0f / 16.0f;

    a = (ushort16)(convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 1]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 2]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 3]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 4]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 5]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 6]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 7]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 8]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 9]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 10]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 11]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 12]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 13]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 14]), convert_ushort(src_p[pixel_index - yScaler * imageWidth - xScaler + 15])
                  );

    b = (ushort16)(convert_ushort(src_p[pixel_index - yScaler * imageWidth]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 1]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 2]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 3]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 4]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 5]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 6]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 7]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 8]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 9]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 10]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 11]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 12]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 13]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + 14]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + 15])
                  );

    c = (ushort16)(convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 1]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 2]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 3]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 4]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 5]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 6]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 7]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 8]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 9]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 10]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 11]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 12]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 13]),
                   convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 14]), convert_ushort(src_p[pixel_index - yScaler * imageWidth + xScaler + 15])
                  );

    d = (ushort16)(convert_ushort(src_p[pixel_index - xScaler]), convert_ushort(src_p[pixel_index - xScaler + 1]),
                   convert_ushort(src_p[pixel_index - xScaler + 2]), convert_ushort(src_p[pixel_index - xScaler + 3]),
                   convert_ushort(src_p[pixel_index - xScaler + 4]), convert_ushort(src_p[pixel_index - xScaler + 5]),
                   convert_ushort(src_p[pixel_index - xScaler + 6]), convert_ushort(src_p[pixel_index - xScaler + 7]),
                   convert_ushort(src_p[pixel_index - xScaler + 8]), convert_ushort(src_p[pixel_index - xScaler + 9]),
                   convert_ushort(src_p[pixel_index - xScaler + 10]), convert_ushort(src_p[pixel_index - xScaler + 11]),
                   convert_ushort(src_p[pixel_index - xScaler + 12]), convert_ushort(src_p[pixel_index - xScaler + 13]),
                   convert_ushort(src_p[pixel_index - xScaler + 14]), convert_ushort(src_p[pixel_index - xScaler + 15])
                  );

    e = (ushort16)(convert_ushort(src_p[pixel_index]), convert_ushort(src_p[pixel_index + 1]),
                   convert_ushort(src_p[pixel_index + 2]), convert_ushort(src_p[pixel_index + 3]),
                   convert_ushort(src_p[pixel_index + 4]), convert_ushort(src_p[pixel_index + 5]),
                   convert_ushort(src_p[pixel_index + 6]), convert_ushort(src_p[pixel_index + 7]),
                   convert_ushort(src_p[pixel_index + 8]), convert_ushort(src_p[pixel_index + 9]),
                   convert_ushort(src_p[pixel_index + 10]), convert_ushort(src_p[pixel_index + 11]),
                   convert_ushort(src_p[pixel_index + 12]), convert_ushort(src_p[pixel_index + 13]),
                   convert_ushort(src_p[pixel_index + 14]), convert_ushort(src_p[pixel_index + 15])
                  );

    f = (ushort16)(convert_ushort(src_p[pixel_index + xScaler]), convert_ushort(src_p[pixel_index + xScaler + 1]),
                   convert_ushort(src_p[pixel_index + xScaler + 2]), convert_ushort(src_p[pixel_index + xScaler + 3]),
                   convert_ushort(src_p[pixel_index + xScaler + 4]), convert_ushort(src_p[pixel_index + xScaler + 5]),
                   convert_ushort(src_p[pixel_index + xScaler + 6]), convert_ushort(src_p[pixel_index + xScaler + 7]),
                   convert_ushort(src_p[pixel_index + xScaler + 8]), convert_ushort(src_p[pixel_index + xScaler + 9]),
                   convert_ushort(src_p[pixel_index + xScaler + 10]), convert_ushort(src_p[pixel_index + xScaler + 11]),
                   convert_ushort(src_p[pixel_index + xScaler + 12]), convert_ushort(src_p[pixel_index + xScaler + 13]),
                   convert_ushort(src_p[pixel_index + xScaler + 14]), convert_ushort(src_p[pixel_index + xScaler + 15])
                  );

    g = (ushort16)(convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 1]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 2]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 3]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 4]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 5]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 6]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 7]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 8]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 9]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 10]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 11]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 12]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 13]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 14]), convert_ushort(src_p[pixel_index + yScaler * imageWidth - xScaler + 15])
                  );

    h = (ushort16)(convert_ushort(src_p[pixel_index + yScaler * imageWidth]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 1]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 2]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 3]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 4]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 5]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 6]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 7]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 8]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 9]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 10]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 11]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 12]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 13]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + 14]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + 15])
                  );

    i = (ushort16)(convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 1]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 2]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 3]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 4]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 5]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 6]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 7]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 8]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 9]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 10]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 11]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 12]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 13]),
                   convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 14]), convert_ushort(src_p[pixel_index + yScaler * imageWidth + xScaler + 15])
                  );

    /*
     { a, b, c } { 1, 2, 1 }
     { d, e, f } { 2, 4, 2 }
     { g, h, i } { 1, 2, 1 }
    */
    ushort16 sum;
    sum = (ushort16)1 * a + (ushort16)2 * b + (ushort16)1 * c +
          (ushort16)2 * d + (ushort16)4 * e + (ushort16)2 * f +
          (ushort16)1 * g + (ushort16)2 * h + (ushort16)1 * i;

    approx = as_uint4(convert_uchar16(((convert_float16(sum) + 0.5f / div) * div)));
    detail = convert_float16(convert_char16(e) - as_char16(approx));

    thold = hardThresh * threshConst[layer - 1];

    detail = (detail < -thold) ? detail + (thold - thold * softThresh) : detail;
    detail = (detail > thold) ? detail - (thold - thold * softThresh) : detail;
    detail = (detail > -thold && detail < thold) ? detail * softThresh : detail;

    __global float16 *details_p = (__global float16 *)(&details[pixel_index]);
    if (layer == 1) {
        (*details_p) = detail;

#if WAVELET_DENOISE_UV
        // copy Y
        luma = vload4(0, src + luma_index0);
        vstore4(luma, 0, dest + luma_index0);
        luma = vload4(0, src + luma_index1);
        vstore4(luma, 0, dest + luma_index1);
#else
        // copy UV
        if (y % 2 == 0) {
            chroma = vload4(0, src + chroma_index + inputUVOffset * (imageWidth / 4));
            vstore4(chroma, 0, dest + chroma_index + outputUVOffset * (imageWidth / 4));
        }
#endif
    } else {
        (*details_p) += detail;
    }

    if (layer < decomLevels) {
#if WAVELET_DENOISE_UV
        int approxOffset = (layer % 2) ? 0 : (inputUVOffset * imageWidth / 4);
        (*(__global uint4*)(approxOut + group_index + approxOffset)) = approx;
#else
        (*(__global uint4*)(approxOut + group_index)) = approx;
#endif
    }
    else
    {
        // Reconstruction
#if WAVELET_DENOISE_UV
        __global uint4* dest_p = (__global uint4*)(&dest[group_index + outputUVOffset * imageWidth / 4]);
        (*dest_p) = as_uint4(convert_uchar16(*details_p + convert_float16(as_uchar16(approx))));
#else
        __global uint4* dest_p = (__global uint4*)(&dest[group_index]);
        (*dest_p) = as_uint4(convert_uchar16(*details_p + convert_float16(as_uchar16(approx))));
#endif
    }
}