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