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