1/* 2 * function: kernel_dark_channel 3 * input_y: Y channel image2d_t as read only 4 * input_uv: UV channel image2d_t as read only 5 * out_dark_channel: dark channel image2d_t as write only 6 * output_r: R channel image2d_t as write only 7 * output_g: G channel image2d_t as write only 8 * output_b: B channel image2d_t as write only 9 * 10 * data_type CL_UNSIGNED_INT16 11 * channel_order CL_RGBA 12 */ 13 14__kernel void kernel_dark_channel ( 15 __read_only image2d_t input_y, __read_only image2d_t input_uv, 16 __write_only image2d_t out_dark_channel, 17 __write_only image2d_t output_r, __write_only image2d_t output_g, __write_only image2d_t output_b) 18{ 19 int pos_x = get_global_id (0); 20 int pos_y = get_global_id (1); 21 sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; 22 float8 y[2]; 23 float8 r, g, b; 24 float8 uv_r, uv_g, uv_b; 25 uint4 ret; 26 int2 pos; 27 28 y[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y * 2))))); 29 y[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y * 2 + 1))))); 30 float8 uv = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_uv, sampler, (int2)(pos_x, pos_y))))) - 128.0f; 31 32 uv_r.even = -0.001f * uv.even + 1.402f * uv.odd; 33 uv_r.odd = uv_r.even; 34 uv_g.even = -0.344f * uv.even - 0.714f * uv.odd; 35 uv_g.odd = uv_g.even; 36 uv_b.even = 1.772f * uv.even + 0.001f * uv.odd; 37 uv_b.odd = uv_b.even; 38 39#pragma unroll 40 for (int i = 0; i < 2; ++i) { 41 r = y[i] + uv_r; 42 g = y[i] + uv_g; 43 b = y[i] + uv_b; 44 r = clamp (r, 0.0f, 255.0f); 45 g = clamp (g, 0.0f, 255.0f); 46 b = clamp (b, 0.0f, 255.0f); 47 48 pos = (int2)(pos_x, 2 * pos_y + i); 49 50 ret = convert_uint4(as_ushort4(convert_uchar8(r))); 51 write_imageui(output_r, pos, ret); 52 ret = convert_uint4(as_ushort4(convert_uchar8(g))); 53 write_imageui(output_g, pos, ret); 54 ret = convert_uint4(as_ushort4(convert_uchar8(b))); 55 write_imageui(output_b, pos, ret); 56 57 r = min (r, g); 58 r = min (r, b); 59 ret = convert_uint4(as_ushort4(convert_uchar8(r))); 60 write_imageui(out_dark_channel, pos, ret); 61 } 62 63} 64 65 66/* 67 * function: kernel_defog_recover 68 * input_dark: dark channel image2d_t as read only 69 * max_v: atmospheric light 70 * input_r: R channel image2d_t as read only 71 * input_g: G channel image2d_t as read only 72 * input_b: B channel image2d_t as read only 73 * output_y: Y channel image2d_t as write only 74 * output_uv: uv channel image2d_t as write only 75 * 76 * data_type CL_UNSIGNED_INT16 77 * channel_order CL_RGBA 78 */ 79 80#define transmit_map_coeff 0.95f 81 82__kernel void kernel_defog_recover ( 83 __read_only image2d_t input_dark, float max_v, float max_r, float max_g, float max_b, 84 __read_only image2d_t input_r, __read_only image2d_t input_g, __read_only image2d_t input_b, 85 __write_only image2d_t out_y, __write_only image2d_t output_uv) 86{ 87 int g_id_x = get_global_id (0); 88 int g_id_y = get_global_id (1); 89 int pos_x = g_id_x; 90 int pos_y = g_id_y * 2; 91 sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; 92 float8 in_r[2], in_g[2], in_b[2]; 93 float8 transmit_map[2]; 94 float8 out_data; 95 96 in_r[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_r, sampler, (int2)(pos_x, pos_y))))); 97 in_r[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_r, sampler, (int2)(pos_x, pos_y + 1))))); 98 in_g[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_g, sampler, (int2)(pos_x, pos_y))))); 99 in_g[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_g, sampler, (int2)(pos_x, pos_y + 1))))); 100 in_b[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_b, sampler, (int2)(pos_x, pos_y))))); 101 in_b[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_b, sampler, (int2)(pos_x, pos_y + 1))))); 102 transmit_map[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x, pos_y))))); 103 transmit_map[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x, pos_y + 1))))); 104 105 transmit_map[0] = 1.0f - transmit_map_coeff * transmit_map[0] / max_v; 106 transmit_map[1] = 1.0f - transmit_map_coeff * transmit_map[1] / max_v; 107 108 transmit_map[0] = max (transmit_map[0], 0.1f); 109 transmit_map[1] = max (transmit_map[1], 0.1f); 110 111 float8 gain = 2.0f; // adjust the brightness temporarily 112 in_r[0] = (max_r + (in_r[0] - max_r) / transmit_map[0]) * gain; 113 in_r[1] = (max_r + (in_r[1] - max_r) / transmit_map[1]) * gain; 114 in_g[0] = (max_g + (in_g[0] - max_g) / transmit_map[0]) * gain; 115 in_g[1] = (max_g + (in_g[1] - max_g) / transmit_map[1]) * gain; 116 in_b[0] = (max_b + (in_b[0] - max_b) / transmit_map[0]) * gain; 117 in_b[1] = (max_b + (in_b[1] - max_b) / transmit_map[1]) * gain; 118 119 out_data = 0.299f * in_r[0] + 0.587f * in_g[0] + 0.114f * in_b[0]; 120 out_data = clamp (out_data, 0.0f, 255.0f); 121 write_imageui(out_y, (int2)(pos_x, pos_y), convert_uint4(as_ushort4(convert_uchar8(out_data)))); 122 out_data = 0.299f * in_r[1] + 0.587f * in_g[1] + 0.114f * in_b[1]; 123 out_data = clamp (out_data, 0.0f, 255.0f); 124 write_imageui(out_y, (int2)(pos_x, pos_y + 1), convert_uint4(as_ushort4(convert_uchar8(out_data)))); 125 126 float4 r, g, b; 127 r = (in_r[0].even + in_r[0].odd + in_r[1].even + in_r[1].odd) * 0.25f; 128 g = (in_g[0].even + in_g[0].odd + in_g[1].even + in_g[1].odd) * 0.25f; 129 b = (in_b[0].even + in_b[0].odd + in_b[1].even + in_b[1].odd) * 0.25f; 130 out_data.even = (-0.169f * r - 0.331f * g + 0.5f * b) + 128.0f; 131 out_data.odd = (0.5f * r - 0.419f * g - 0.081f * b) + 128.0f; 132 out_data = clamp (out_data, 0.0f, 255.0f); 133 write_imageui(output_uv, (int2)(g_id_x, g_id_y), convert_uint4(as_ushort4(convert_uchar8(out_data)))); 134} 135 136