cl_kernel/kernel_defog_dcp.cl

/*
 * function:    kernel_dark_channel
 * input_y:     Y channel image2d_t as read only
 * input_uv:   UV channel image2d_t as read only
 * out_dark_channel: dark channel image2d_t as write only
 * output_r:    R channel image2d_t as write only
 * output_g:   G channel image2d_t as write only
 * output_b:   B channel image2d_t as write only
 *
 * data_type CL_UNSIGNED_INT16
 * channel_order CL_RGBA
 */

__kernel void kernel_dark_channel (
    __read_only image2d_t input_y, __read_only image2d_t input_uv,
    __write_only image2d_t out_dark_channel,
    __write_only image2d_t output_r, __write_only image2d_t output_g, __write_only image2d_t output_b)
{
    int pos_x = get_global_id (0);
    int pos_y = get_global_id (1);
    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    float8 y[2];
    float8 r, g, b;
    float8 uv_r, uv_g, uv_b;
    uint4 ret;
    int2 pos;

    y[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y * 2)))));
    y[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y * 2 + 1)))));
    float8 uv = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_uv, sampler, (int2)(pos_x, pos_y))))) - 128.0f;

    uv_r.even = -0.001f * uv.even + 1.402f * uv.odd;
    uv_r.odd = uv_r.even;
    uv_g.even = -0.344f * uv.even - 0.714f * uv.odd;
    uv_g.odd = uv_g.even;
    uv_b.even = 1.772f * uv.even + 0.001f * uv.odd;
    uv_b.odd = uv_b.even;

#pragma unroll
    for (int i = 0; i < 2; ++i) {
        r = y[i] + uv_r;
        g = y[i] + uv_g;
        b = y[i] + uv_b;
        r = clamp (r, 0.0f, 255.0f);
        g = clamp (g, 0.0f, 255.0f);
        b = clamp (b, 0.0f, 255.0f);

        pos = (int2)(pos_x, 2 * pos_y + i);

        ret = convert_uint4(as_ushort4(convert_uchar8(r)));
        write_imageui(output_r, pos, ret);
        ret = convert_uint4(as_ushort4(convert_uchar8(g)));
        write_imageui(output_g, pos, ret);
        ret = convert_uint4(as_ushort4(convert_uchar8(b)));
        write_imageui(output_b, pos, ret);

        r = min (r, g);
        r = min (r, b);
        ret = convert_uint4(as_ushort4(convert_uchar8(r)));
        write_imageui(out_dark_channel, pos, ret);
    }

}


/*
 * function:    kernel_defog_recover
 * input_dark: dark channel image2d_t as read only
 * max_v:       atmospheric light
 * input_r:      R channel image2d_t as read only
 * input_g:     G channel image2d_t as read only
 * input_b:     B channel image2d_t as read only
 * output_y:   Y channel image2d_t as write only
 * output_uv: uv channel image2d_t as write only
 *
 * data_type        CL_UNSIGNED_INT16
 * channel_order  CL_RGBA
 */

#define transmit_map_coeff 0.95f

__kernel void kernel_defog_recover (
    __read_only image2d_t input_dark, float max_v, float max_r, float max_g, float max_b,
    __read_only image2d_t input_r, __read_only image2d_t input_g, __read_only image2d_t input_b,
    __write_only image2d_t out_y, __write_only image2d_t output_uv)
{
    int g_id_x = get_global_id (0);
    int g_id_y = get_global_id (1);
    int pos_x = g_id_x;
    int pos_y = g_id_y * 2;
    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    float8 in_r[2], in_g[2], in_b[2];
    float8 transmit_map[2];
    float8 out_data;

    in_r[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_r, sampler, (int2)(pos_x, pos_y)))));
    in_r[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_r, sampler, (int2)(pos_x, pos_y + 1)))));
    in_g[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_g, sampler, (int2)(pos_x, pos_y)))));
    in_g[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_g, sampler, (int2)(pos_x, pos_y + 1)))));
    in_b[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_b, sampler, (int2)(pos_x, pos_y)))));
    in_b[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_b, sampler, (int2)(pos_x, pos_y + 1)))));
    transmit_map[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x, pos_y)))));
    transmit_map[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x, pos_y + 1)))));

    transmit_map[0] = 1.0f - transmit_map_coeff * transmit_map[0] / max_v;
    transmit_map[1] = 1.0f - transmit_map_coeff * transmit_map[1] / max_v;

    transmit_map[0] = max (transmit_map[0], 0.1f);
    transmit_map[1] = max (transmit_map[1], 0.1f);

    float8 gain = 2.0f;  // adjust the brightness temporarily
    in_r[0] = (max_r + (in_r[0] - max_r) / transmit_map[0]) * gain;
    in_r[1] = (max_r + (in_r[1] - max_r) / transmit_map[1]) * gain;
    in_g[0] = (max_g + (in_g[0] - max_g) / transmit_map[0]) * gain;
    in_g[1] = (max_g + (in_g[1] - max_g) / transmit_map[1]) * gain;
    in_b[0] = (max_b + (in_b[0] - max_b) / transmit_map[0]) * gain;
    in_b[1] = (max_b + (in_b[1] - max_b) / transmit_map[1]) * gain;

    out_data = 0.299f * in_r[0] + 0.587f * in_g[0] + 0.114f * in_b[0];
    out_data = clamp (out_data, 0.0f, 255.0f);
    write_imageui(out_y, (int2)(pos_x, pos_y), convert_uint4(as_ushort4(convert_uchar8(out_data))));
    out_data = 0.299f * in_r[1] + 0.587f * in_g[1] + 0.114f * in_b[1];
    out_data = clamp (out_data, 0.0f, 255.0f);
    write_imageui(out_y, (int2)(pos_x, pos_y + 1), convert_uint4(as_ushort4(convert_uchar8(out_data))));

    float4 r, g, b;
    r = (in_r[0].even + in_r[0].odd + in_r[1].even + in_r[1].odd) * 0.25f;
    g = (in_g[0].even + in_g[0].odd + in_g[1].even + in_g[1].odd) * 0.25f;
    b = (in_b[0].even + in_b[0].odd + in_b[1].even + in_b[1].odd) * 0.25f;
    out_data.even = (-0.169f * r - 0.331f * g + 0.5f * b) + 128.0f;
    out_data.odd = (0.5f * r - 0.419f * g - 0.081f * b) + 128.0f;
    out_data = clamp (out_data, 0.0f, 255.0f);
    write_imageui(output_uv, (int2)(g_id_x, g_id_y), convert_uint4(as_ushort4(convert_uchar8(out_data))));
}