| /* |
| * function: kernel_bayer_copy |
| * sample code of default kernel arguments |
| * input: image2d_t as read only |
| * output: image2d_t as write only |
| */ |
| |
| //#define ENABLE_IMAGE_2D_INPUT 0 |
| |
| #ifndef STATS_BITS |
| #define STATS_BITS 8 |
| #endif |
| |
| /* |
| * GROUP_PIXEL_X_SIZE = 2 * GROUP_CELL_X_SIZE |
| * GROUP_PIXEL_Y_SIZE = 2 * GROUP_CELL_Y_SIZE |
| */ |
| |
| #define GROUP_CELL_X_SIZE 64 |
| #define GROUP_CELL_Y_SIZE 4 |
| |
| //float4; 16 |
| #define SLM_X_SIZE (GROUP_CELL_X_SIZE / 4) |
| #define SLM_Y_SIZE GROUP_CELL_Y_SIZE |
| |
| #define STATS_3A_CELL_X_SIZE 8 |
| #define STATS_3A_CELL_Y_SIZE GROUP_CELL_Y_SIZE |
| |
| typedef struct { |
| float level_gr; /* Black level for GR pixels */ |
| float level_r; /* Black level for R pixels */ |
| float level_b; /* Black level for B pixels */ |
| float level_gb; /* Black level for GB pixels */ |
| uint color_bits; |
| } CLBLCConfig; |
| |
| |
| typedef struct |
| { |
| float r_gain; |
| float gr_gain; |
| float gb_gain; |
| float b_gain; |
| } CLWBConfig; |
| |
| inline int slm_pos (const int x, const int y) |
| { |
| return mad24 (y, SLM_X_SIZE, x); |
| } |
| |
| inline void gamma_correct(float8 *in_out, __global float *table) |
| { |
| in_out->s0 = table[clamp(convert_int(in_out->s0 * 255.0f), 0, 255)]; |
| in_out->s1 = table[clamp(convert_int(in_out->s1 * 255.0f), 0, 255)]; |
| in_out->s2 = table[clamp(convert_int(in_out->s2 * 255.0f), 0, 255)]; |
| in_out->s3 = table[clamp(convert_int(in_out->s3 * 255.0f), 0, 255)]; |
| in_out->s4 = table[clamp(convert_int(in_out->s4 * 255.0f), 0, 255)]; |
| in_out->s5 = table[clamp(convert_int(in_out->s5 * 255.0f), 0, 255)]; |
| in_out->s6 = table[clamp(convert_int(in_out->s6 * 255.0f), 0, 255)]; |
| in_out->s7 = table[clamp(convert_int(in_out->s7 * 255.0f), 0, 255)]; |
| } |
| |
| inline float avg_float8 (float8 data) |
| { |
| return (data.s0 + data.s1 + data.s2 + data.s3 + data.s4 + data.s5 + data.s6 + data.s7) * 0.125f; |
| } |
| |
| inline void stats_3a_calculate ( |
| __local float4 * slm_gr, |
| __local float4 * slm_r, |
| __local float4 * slm_b, |
| __local float4 * slm_gb, |
| __global ushort8 * stats_output, |
| CLWBConfig *wb_config) |
| { |
| const int group_x_size = get_num_groups (0); |
| const int group_id_x = get_group_id (0); |
| const int group_id_y = get_group_id (1); |
| |
| const int l_id_x = get_local_id (0); |
| const int l_id_y = get_local_id (1); |
| const int l_size_x = get_local_size (0); |
| const int stats_float4_x_count = STATS_3A_CELL_X_SIZE / 4; |
| int count = stats_float4_x_count * STATS_3A_CELL_Y_SIZE / 4; |
| |
| int index = mad24 (l_id_y, l_size_x, l_id_x); |
| int index_x = index % SLM_X_SIZE; |
| int index_y = index / SLM_X_SIZE; |
| |
| if (mad24 (index_y, stats_float4_x_count, index_x % stats_float4_x_count) < count) { |
| int pitch_count = count / stats_float4_x_count * SLM_X_SIZE; |
| int index1 = index + pitch_count; |
| int index2 = index1 + pitch_count; |
| int index3 = index2 + pitch_count; |
| slm_gr[index] = (slm_gr[index] + slm_gr[index1] + slm_gr[index2] + slm_gr[index3]) * 0.25f; |
| slm_r[index] = (slm_r[index] + slm_r[index1] + slm_r[index2] + slm_r[index3]) * 0.25f; |
| slm_b[index] = (slm_b[index] + slm_b[index1] + slm_b[index2] + slm_b[index3]) * 0.25f; |
| slm_gb[index] = (slm_gb[index] + slm_gb[index1] + slm_gb[index2] + slm_gb[index3]) * 0.25f; |
| } |
| barrier (CLK_LOCAL_MEM_FENCE); |
| |
| if (index < SLM_X_SIZE / 2) { |
| float result_gr, result_r, result_b, result_gb, avg_y; |
| float8 tmp; |
| tmp = ((__local float8*)slm_gr)[index]; |
| result_gr = avg_float8 (tmp); |
| |
| tmp = ((__local float8*)slm_r)[index]; |
| result_r = avg_float8 (tmp); |
| |
| tmp = ((__local float8*)slm_b)[index]; |
| result_b = avg_float8 (tmp); |
| |
| tmp = ((__local float8*)slm_gb)[index]; |
| result_gb = avg_float8 (tmp); |
| |
| int out_index = mad24 (mad24 (group_id_y, group_x_size, group_id_x), |
| (GROUP_CELL_X_SIZE / STATS_3A_CELL_X_SIZE) * (GROUP_CELL_Y_SIZE / STATS_3A_CELL_Y_SIZE), |
| index); |
| |
| #if STATS_BITS==8 |
| avg_y = mad ((result_gr * wb_config->gr_gain + result_gb * wb_config->gb_gain), 74.843f, |
| mad (result_r * wb_config->r_gain, 76.245f, result_b * 29.070f)); |
| |
| //ushort avg_y; avg_r; avg_gr; avg_gb; avg_b; valid_wb_count; f_value1; f_value2; |
| stats_output[out_index] = (ushort8) ( |
| convert_ushort (convert_uchar_sat (avg_y)), |
| convert_ushort (convert_uchar_sat (result_r * 255.0f)), |
| convert_ushort (convert_uchar_sat (result_gr * 255.0f)), |
| convert_ushort (convert_uchar_sat (result_gb * 255.0f)), |
| convert_ushort (convert_uchar_sat (result_b * 255.0f)), |
| STATS_3A_CELL_X_SIZE * STATS_3A_CELL_Y_SIZE, |
| 0, |
| 0); |
| #elif STATS_BITS==12 |
| avg_y = mad ((result_gr * wb_config->gr_gain + result_gb * wb_config->gb_gain), 1201.883f, |
| mad (result_r * wb_config->r_gain, 1224.405f, result_b * 466.830f)); |
| |
| stats_output[out_index] = (ushort8) ( |
| convert_ushort (clamp (avg_y, 0.0f, 4095.0f)), |
| convert_ushort (clamp (result_r * 4096.0f, 0.0f, 4095.0f)), |
| convert_ushort (clamp (result_gr * 4096.0f, 0.0f, 4095.0f)), |
| convert_ushort (clamp (result_gb * 4096.0f, 0.0f, 4095.0f)), |
| convert_ushort (clamp (result_b * 4096.0f, 0.0f, 4095.0f)), |
| STATS_3A_CELL_X_SIZE * STATS_3A_CELL_Y_SIZE, |
| 0, |
| 0); |
| #else |
| printf ("kernel 3a-stats error, wrong bit depth:%d\n", STATS_BITS); |
| #endif |
| } |
| } |
| |
| |
| __kernel void kernel_bayer_basic ( |
| #if ENABLE_IMAGE_2D_INPUT |
| __read_only image2d_t input, |
| #else |
| __global const ushort8 *input, |
| #endif |
| uint input_aligned_width, |
| __write_only image2d_t output, |
| uint out_height, |
| CLBLCConfig blc_config, |
| CLWBConfig wb_config, |
| __global float *gamma_table, |
| __global ushort8 *stats_output |
| ) |
| { |
| int g_x = get_global_id (0); |
| int g_y = get_global_id (1); |
| |
| const int l_x = get_local_id (0); |
| const int l_y = get_local_id (1); |
| const int l_x_size = get_local_size (0); |
| const int l_y_size = get_local_size (1); |
| const int group_id_x = get_group_id (0); |
| const int group_id_y = get_group_id (1); |
| |
| sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; |
| |
| int index = mad24 (l_y, l_x_size, l_x); |
| int x_cell_start = (GROUP_CELL_X_SIZE / 4) * group_id_x; |
| int y_cell_start = GROUP_CELL_Y_SIZE * group_id_y; |
| int x, y; |
| |
| float blc_multiplier = (float)(1 << (16 - blc_config.color_bits)); |
| |
| __local float4 slm_gr[SLM_X_SIZE * SLM_Y_SIZE], slm_r[SLM_X_SIZE * SLM_Y_SIZE], slm_b[SLM_X_SIZE * SLM_Y_SIZE], slm_gb[SLM_X_SIZE * SLM_Y_SIZE]; |
| |
| for (; index < SLM_X_SIZE * SLM_Y_SIZE; index += l_x_size * l_y_size) { |
| float8 line1; |
| float8 line2; |
| |
| x = index % SLM_X_SIZE + x_cell_start; |
| y = index / SLM_X_SIZE + y_cell_start; |
| |
| #if ENABLE_IMAGE_2D_INPUT |
| line1 = convert_float8 (as_ushort8 (read_imageui(input, sampler, (int2)(x, y * 2)))) / 65536.0f; |
| line2 = convert_float8 (as_ushort8 (read_imageui(input, sampler, (int2)(x, y * 2 + 1)))) / 65536.0f; |
| #else |
| line1 = convert_float8 (input [y * 2 * input_aligned_width + x]) / 65536.0f; |
| line2 = convert_float8 (input [(y * 2 + 1) * input_aligned_width + x]) / 65536.0f; |
| #endif |
| |
| float4 gr = mad (line1.even, blc_multiplier, - blc_config.level_gr); |
| float4 r = mad (line1.odd, blc_multiplier, - blc_config.level_r); |
| float4 b = mad (line2.even, blc_multiplier, - blc_config.level_b); |
| float4 gb = mad (line2.odd, blc_multiplier, - blc_config.level_gb); |
| |
| slm_gr[index] = gr; |
| slm_r[index] = r; |
| slm_b[index] = b; |
| slm_gb[index] = gb; |
| } |
| barrier(CLK_LOCAL_MEM_FENCE); |
| |
| float8 data_gr, data_r, data_b, data_gb; |
| index = mad24 (l_y, l_x_size, l_x); |
| x = mad24 (GROUP_CELL_X_SIZE / 8, group_id_x, index % (SLM_X_SIZE / 2)); |
| y = mad24 (GROUP_CELL_Y_SIZE, group_id_y, index / (SLM_X_SIZE / 2)); |
| |
| data_gr = ((__local float8*)slm_gr)[index]; |
| data_gr = data_gr * wb_config.gr_gain; |
| |
| data_r = ((__local float8*)slm_r)[index]; |
| data_r = data_r * wb_config.r_gain; |
| |
| data_b = ((__local float8*)slm_b)[index]; |
| data_b = data_b * wb_config.b_gain; |
| |
| data_gb = ((__local float8*)slm_gb)[index]; |
| data_gb = data_gb * wb_config.gb_gain; |
| |
| #if ENABLE_GAMMA |
| gamma_correct (&data_gr, gamma_table); |
| gamma_correct (&data_r, gamma_table); |
| gamma_correct (&data_b, gamma_table); |
| gamma_correct (&data_gb, gamma_table); |
| #endif |
| |
| #if 0 |
| if (x % 16 == 0 && y % 16 == 0) { |
| uint8 value = convert_uint8(convert_uchar8_sat(data_gr * 255.0f)); |
| printf ("(x:%d, y:%d) (blc.bit:%d, level:%d) (wb.gr:%f)=> (%d, %d, %d, %d, %d, %d, %d, %d)\n", |
| x * 8, y, |
| blc_config.color_bits, convert_uint(blc_config.level_gr * 255.0f), |
| wb_config.gr_gain, |
| value.s0, value.s1, value.s2, value.s3, value.s4, value.s5, value.s6, value.s7); |
| } |
| #endif |
| |
| write_imageui (output, (int2)(x, y), as_uint4 (convert_ushort8 (data_gr * 65536.0f))); |
| write_imageui (output, (int2)(x, y + out_height), as_uint4 (convert_ushort8 (data_r * 65536.0f))); |
| write_imageui (output, (int2)(x, y + out_height * 2), as_uint4 (convert_ushort8 (data_b * 65536.0f))); |
| write_imageui (output, (int2)(x, y + out_height * 3), as_uint4 (convert_ushort8 (data_gb * 65536.0f))); |
| |
| stats_3a_calculate (slm_gr, slm_r, slm_b, slm_gb, stats_output, &wb_config); |
| } |
| |
