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