/*
* 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))));
}