#ifdef MNN_SUPPORT_FP16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#endif
#define GLOBAL_SIZE_3_DIMS \
    __private const int global_size_dim0, __private const int global_size_dim1, __private const int global_size_dim2,
#define DEAL_NON_UNIFORM_DIM3(input1, input2, input3)                                             \
    if (input1 >= global_size_dim0 || input2 >= global_size_dim1 || input3 >= global_size_dim2) { \
        return;                                                                                   \
    }

__kernel void pooling_c4_c4(GLOBAL_SIZE_3_DIMS __global const FLOAT *input,
                      __private const int2 input_shape,
                      __private const int2 output_shape,
                      __private const int2 pad_shape,
                      __global FLOAT *output,
                      __global FLOAT *rediceOutput,
                      __private const int channel,
                      __private const int batch,
                      __private const int in_channel_block,
                      __private const int out_channel_block,
                      __private const int input_pad_left,
                      __private const int input_pad_right, 
                      __private const int output_pad_left,
                      __private const int output_pad_right) {
                          
    const int ow_idx   = get_global_id(0);
    const int b_oh_idx = get_global_id(1);
    const int c_idx    = get_global_id(2);

    DEAL_NON_UNIFORM_DIM3(ow_idx, b_oh_idx, c_idx);
    
    const int b_idx = b_oh_idx / output_shape.x;
    const int oh_idx = b_oh_idx % output_shape.x;
    const int iw_start = mad24(ow_idx, STRIDE_X, -pad_shape.y);
    const int ih_start = mad24(oh_idx, STRIDE_Y, -pad_shape.x);
    
    #ifdef POOL_AVG
    COMPUTE_FLOAT4 result = (COMPUTE_FLOAT4)(0);
    const int inp_offset = (((b_idx+c_idx*batch)*input_shape.x+ih_start)*input_shape.y+iw_start+input_pad_left)*4;
#ifdef COUNT_INCLUDE_PADDING
    int total_count = (min(ih_start + KERNEL_Y, input_shape.x + pad_shape.x) - ih_start) * (min(iw_start + KERNEL_X, input_shape.y + pad_shape.y) - iw_start);
#else
    int total_count = 0;
#endif
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }
        for(int kw=0; kw<KERNEL_X; kw++) {
            int iw_cur = iw_start + kw;
            if(iw_cur < 0 || iw_cur >= input_shape.y) {
                continue;
            }
            COMPUTE_FLOAT4 inp_data = CONVERT_COMPUTE_FLOAT4(vload4(0, input+inp_offset+(kh*input_shape.y+kw)*4));
            result += inp_data;
#ifndef COUNT_INCLUDE_PADDING
            total_count++;
#endif
        }
    }
    result = result / (COMPUTE_FLOAT4)(1.0*total_count);
    #else
    COMPUTE_FLOAT4 result = (COMPUTE_FLOAT4)(-FLT_MAX);
    #if RETURN_REDICE
    int4 redice = (int4)0;
    #endif
    const int inp_offset = (((b_idx+c_idx*batch)*input_shape.x+ih_start)*input_shape.y+iw_start+input_pad_left)*4;
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }
        for(int kw=0; kw<KERNEL_X; kw++) {
            int iw_cur = iw_start + kw;
            if(iw_cur < 0 || iw_cur >= input_shape.y) {
                continue;
            }
            COMPUTE_FLOAT4 inp_data = CONVERT_COMPUTE_FLOAT4(vload4(0, input+inp_offset+(kh*input_shape.y+kw)*4));
            #if RETURN_REDICE
            redice = inp_data > result ? (int4)((ih_start + kh) * input_shape.y + iw_start + kw) : redice;
            #endif
            result = fmax(result, inp_data);
        }
    }
    #endif
    
    const int out_offset = (((b_idx + c_idx*batch)*output_shape.x + oh_idx)* output_shape.y + ow_idx + output_pad_left)*4;
    vstore4(CONVERT_FLOAT4(result), 0, output+out_offset);
    #if RETURN_REDICE
    vstore4(CONVERT_FLOAT4(redice),  0, rediceOutput+(((b_idx + c_idx*batch)*output_shape.x + oh_idx)* output_shape.y + ow_idx)*4);
    #endif
}

__kernel void pooling_c4_c16(GLOBAL_SIZE_3_DIMS __global const FLOAT *input,
                      __private const int2 input_shape,
                      __private const int2 output_shape,
                      __private const int2 pad_shape,
                      __global FLOAT *output,
                      __global FLOAT *rediceOutput,
                      __private const int channel,
                      __private const int batch,
                      __private const int in_channel_block,
                      __private const int out_channel_block,
                      __private const int input_pad_left, 
                      __private const int input_pad_right, 
                      __private const int output_pad_left,
                      __private const int output_pad_right) {
                          
    const int ow_idx   = get_global_id(0);
    const int b_oh_idx = get_global_id(1);
    const int c_idx    = get_global_id(2);

    DEAL_NON_UNIFORM_DIM3(ow_idx, b_oh_idx, c_idx);
    
    const int b_idx = b_oh_idx / output_shape.x;
    const int oh_idx = b_oh_idx % output_shape.x;
    const int iw_start = mad24(ow_idx, STRIDE_X, -pad_shape.y);
    const int ih_start = mad24(oh_idx, STRIDE_Y, -pad_shape.x);
    const int dst_width = output_shape.y + output_pad_left + output_pad_right;
    
    #ifdef POOL_AVG
    COMPUTE_FLOAT4 result = (COMPUTE_FLOAT4)(0);
    const int inp_offset = (((b_idx+c_idx*batch)*input_shape.x+ih_start)*input_shape.y+iw_start+input_pad_left)*4;
 #ifdef COUNT_INCLUDE_PADDING
    int total_count = (min(ih_start + KERNEL_Y, input_shape.x + pad_shape.x) - ih_start) * (min(iw_start + KERNEL_X, input_shape.y + pad_shape.y) - iw_start);
#else
    int total_count = 0;
#endif
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }
        for(int kw=0; kw<KERNEL_X; kw++) {
            int iw_cur = iw_start + kw;
            if(iw_cur < 0 || iw_cur >= input_shape.y) {
                continue;
            }
            COMPUTE_FLOAT4 inp_data = CONVERT_COMPUTE_FLOAT4(vload4(0, input+inp_offset+(kh*input_shape.y+kw)*4));
            result += inp_data;
#ifndef COUNT_INCLUDE_PADDING
            total_count++;
#endif
        }
    }
    result = result / (COMPUTE_FLOAT4)(1.0*total_count);
    #else
    COMPUTE_FLOAT4 result = (COMPUTE_FLOAT4)(-FLT_MAX);
    #if RETURN_REDICE
    int4 redice = (int4)0;
    #endif
    const int inp_offset = (((b_idx+c_idx*batch)*input_shape.x+ih_start)*input_shape.y+iw_start+input_pad_left)*4;
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }
        for(int kw=0; kw<KERNEL_X; kw++) {
            int iw_cur = iw_start + kw;
            if(iw_cur < 0 || iw_cur >= input_shape.y) {
                continue;
            }
            COMPUTE_FLOAT4 inp_data = CONVERT_COMPUTE_FLOAT4(vload4(0, input+inp_offset+(kh*input_shape.y+kw)*4));
            #if RETURN_REDICE
            redice = inp_data > result ? (int4)((ih_start + kh) * input_shape.y + iw_start + kw) : redice;
            #endif
            result = fmax(result, inp_data);
        }
    }
    #endif

    const int c_left = (c_idx % 4) * 4;
    const int out_offset = (((b_idx*out_channel_block + c_idx/4)*output_shape.x + oh_idx)* dst_width + ow_idx + output_pad_left)*16 + c_left;
    vstore4(CONVERT_FLOAT4(result), 0, output+out_offset);
    #if RETURN_REDICE
    vstore4(CONVERT_FLOAT4(redice),  0, rediceOutput+(((b_idx*out_channel_block + c_idx)*output_shape.x + oh_idx)* output_shape.y + ow_idx)*4);
    #endif
    if(ow_idx == 0){
        int pad_offset = (((b_idx*out_channel_block + c_idx/4)*output_shape.x + oh_idx)* dst_width + 0)*16 + c_left;
        for(int i = 0; i < output_pad_left; ++i){
            vstore4((FLOAT4)0, 0, output + pad_offset + i * 16);
        }
        pad_offset += (output_shape.y + output_pad_left) * 16;
        for(int i = 0; i < output_pad_right; ++i){
            vstore4((FLOAT4)0, 0, output + pad_offset + i * 16);
        }
    }
}

__attribute__((intel_reqd_sub_group_size(16)))
__kernel void pooling_c16_c16(GLOBAL_SIZE_3_DIMS __global const FLOAT *input,
                      __private const int2 input_shape,
                      __private const int2 output_shape,
                      __private const int2 pad_shape,
                      __global FLOAT *output,
                      __global FLOAT *rediceOutput,
                      __private const int channel,
                      __private const int batch,
                      __private const int in_channel_block,
                      __private const int out_channel_block,
                      __private const int input_pad_left, 
                      __private const int input_pad_right, 
                      __private const int output_pad_left,
                      __private const int output_pad_right) {
                          
    const int ow_idx   = get_global_id(1) << 3;
    const int b_oh_idx = get_global_id(2);
    const int c_idx    = get_group_id(0);
    const int sglid = get_sub_group_local_id();
    
    const int b_idx = b_oh_idx / output_shape.x;
    const int oh_idx = b_oh_idx % output_shape.x;
    const int iw_start = mad24(ow_idx, STRIDE_X, -pad_shape.y);
    const int ih_start = mad24(oh_idx, STRIDE_Y, -pad_shape.x);
    const int src_width = input_shape.y + input_pad_left + input_pad_right;
    const int dst_width = output_shape.y + output_pad_left + output_pad_right;

#ifdef POOL_AVG
    COMPUTE_FLOAT8 result = (COMPUTE_FLOAT8)(0);
    COMPUTE_FLOAT8 w_start = (COMPUTE_FLOAT8)(iw_start, iw_start + STRIDE_X, iw_start + STRIDE_X * 2, iw_start + STRIDE_X * 3, iw_start + STRIDE_X * 4, iw_start + STRIDE_X * 5, iw_start + STRIDE_X * 6, iw_start + STRIDE_X * 7);
#ifdef COUNT_INCLUDE_PADDING
    COMPUTE_FLOAT8 w_size      = fmin(w_start + KERNEL_X, input_shape.y + pad_shape.y) - w_start;
    COMPUTE_FLOAT8 total_count = (COMPUTE_FLOAT8)(min(ih_start + KERNEL_Y, input_shape.x + pad_shape.x) - ih_start) * w_size;
#else
    w_start = fmax(w_start, (COMPUTE_FLOAT8)0);
    COMPUTE_FLOAT8 w_end = fmin(w_start + KERNEL_X, (COMPUTE_FLOAT8)input_shape.y);
    float h_start = fmax((float)ih_start, 0);
    float h_end = fmin(h_start + KERNEL_Y, (float)input_shape.x);
    COMPUTE_FLOAT8 total_count = (w_end - w_start) * (COMPUTE_FLOAT8)(h_end - h_start);
#endif
#else
    COMPUTE_FLOAT8 result = (COMPUTE_FLOAT8)(-FLT_MAX);
#if RETURN_REDICE
    int8 redice = (int8)0;
#endif
#endif
    const int inp_offset = mul24(mad24(mad24(mad24(b_idx,in_channel_block,c_idx),input_shape.x,ih_start),src_width,iw_start+input_pad_left),16);
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }

        FLOAT line_cache[INPUT_LINE_SIZE];
        for (int i = 0; i < INPUT_LINE_SIZE; i++) {
            if ((iw_start + i) >= 0 && (iw_start + i) < input_shape.y){
#ifdef MNN_SUPPORT_FP16
                line_cache[i] = as_half(intel_sub_group_block_read_us((__global ushort*)(input + inp_offset + mul24(mad24(kh,src_width,i),16))));
#else
                line_cache[i] = as_float(intel_sub_group_block_read((__global uint*)(input + inp_offset + mul24(mad24(kh,src_width,i),16))));
#endif
            } else{
#ifdef POOL_AVG
                line_cache[i] = 0;
#else
                line_cache[i] = (COMPUTE_FLOAT)(-FLT_MAX);
#endif
            }
        }


        for(int kw=0; kw<KERNEL_X; kw++) {
            COMPUTE_FLOAT8 src;
            __attribute__((opencl_unroll_hint(8)))
            for (int i = 0; i < 8; i++) {
                src[i] = line_cache[kw + STRIDE_X*i];
            }
#ifdef POOL_AVG
            result += src;
#else
#if RETURN_REDICE
            redice = src > result ? (int8)((ih_start + kh) * input_shape.y + iw_start + kw) : redice;
#endif
            result = fmax(result, src);
#endif
        }
    }
#ifdef POOL_AVG
    result = result / total_count;
#endif


    if(ow_idx == 0){
        int pad_offset = (((b_idx*out_channel_block + c_idx)*output_shape.x + oh_idx)* dst_width + 0)*16 + sglid;
        for(int i = 0; i < output_pad_left; ++i){
            output[pad_offset+i*16] = 0;
        }
        pad_offset += (output_shape.y + output_pad_left) * 16;
        for(int i = 0; i < output_pad_right; ++i){
            output[pad_offset+i*16] = 0;
        }
    }
    
    const int out_offset = (((b_idx*out_channel_block + c_idx)*output_shape.x + oh_idx)* dst_width + ow_idx + output_pad_left)*16;
#if OUTPUT_LEFTOVERS
    if ((c_idx+1)*16 >= channel) {
        for (int i = 0; i < 8; i++) {
            if ((c_idx*16 + sglid < channel) && (ow_idx + i) < output_shape.y)
                output[out_offset + i * 16 + sglid] = result[i];
        }
    }
    else
#endif  
    {
        if (ow_idx + 8 <= output_shape.y) {
#ifdef MNN_SUPPORT_FP16
            intel_sub_group_block_write_us8((__global ushort*)(output + out_offset), as_ushort8(CONVERT_FLOAT8(result)));
#else
            intel_sub_group_block_write8((__global uint*)(output + out_offset), as_uint8(CONVERT_FLOAT8(result)));
#endif
        }else{
            for (int i = 0; i < output_shape.y % 8; i++) {
                output[out_offset + i * 16 + sglid] = result[i];
            }
        }
    }
#ifdef RETURN_REDICE
    const uint lid_x = sglid % 4;
    const uint lid_y = sglid / 4;
    
    const int width_height = output_shape.y * output_shape.x * 4;
    const int redice_offset = (((b_idx*out_channel_block + c_idx * 4)*output_shape.x + oh_idx)* output_shape.y + ow_idx)*4;
#if OUTPUT_LEFTOVERS
    if ((c_idx+1)*16 >= channel) {
        for (int i = 0; i < 8; i++) {
            if ((c_idx*16 + lid_y * 4 + lid_x < channel) && (ow_idx + i) < output_shape.y)
                rediceOutput[redice_offset + lid_y * width_height + i * 4 + lid_x] = redice[i];
        }
    }
    else
#endif
    {
        for (int i = 0; i < 8 && (ow_idx + i) < output_shape.y; i++) {
            rediceOutput[redice_offset + lid_y * width_height + i * 4 + lid_x] = redice[i];
        }
    }
#endif
}

__attribute__((intel_reqd_sub_group_size(16)))
__kernel void pooling_c16_c4(GLOBAL_SIZE_3_DIMS __global const FLOAT *input,
                      __private const int2 input_shape,
                      __private const int2 output_shape,
                      __private const int2 pad_shape,
                      __global FLOAT *output,
                      __global FLOAT *rediceOutput,
                      __private const int channel,
                      __private const int batch,
                      __private const int in_channel_block,
                      __private const int out_channel_block,
                      __private const int input_pad_left, 
                      __private const int input_pad_right, 
                      __private const int output_pad_left,
                      __private const int output_pad_right) {
                          
    const int ow_idx   = get_global_id(1) << 3;
    const int b_oh_idx = get_global_id(2);
    const int c_idx    = get_group_id(0);
    const int sglid = get_sub_group_local_id();
    
    const int b_idx = b_oh_idx / output_shape.x;
    const int oh_idx = b_oh_idx % output_shape.x;
    const int iw_start = mad24(ow_idx, STRIDE_X, -pad_shape.y);
    const int ih_start = mad24(oh_idx, STRIDE_Y, -pad_shape.x);
    const int src_width = input_shape.y + input_pad_left + input_pad_right;

#ifdef POOL_AVG
    COMPUTE_FLOAT8 result = (COMPUTE_FLOAT8)(0);
    COMPUTE_FLOAT8 w_start = (COMPUTE_FLOAT8)(iw_start, iw_start + STRIDE_X, iw_start + STRIDE_X * 2, iw_start + STRIDE_X * 3, iw_start + STRIDE_X * 4, iw_start + STRIDE_X * 5, iw_start + STRIDE_X * 6, iw_start + STRIDE_X * 7);
#ifdef COUNT_INCLUDE_PADDING
    COMPUTE_FLOAT8 w_size      = fmin(w_start + KERNEL_X, input_shape.y + pad_shape.y) - w_start;
    COMPUTE_FLOAT8 total_count = (COMPUTE_FLOAT8)(min(ih_start + KERNEL_Y, input_shape.x + pad_shape.x) - ih_start) * w_size;
#else
    w_start = fmax(w_start, (COMPUTE_FLOAT8)0);
    COMPUTE_FLOAT8 w_end = fmin(w_start + KERNEL_X, (COMPUTE_FLOAT8)input_shape.y);
    float h_start = fmax((float)ih_start, 0);
    float h_end = fmin(h_start + KERNEL_Y, (float)input_shape.x);
    COMPUTE_FLOAT8 total_count = (w_end - w_start) * (COMPUTE_FLOAT8)(h_end - h_start);
#endif
#else
    COMPUTE_FLOAT8 result = (COMPUTE_FLOAT8)(-FLT_MAX);
#if RETURN_REDICE
    int8 redice = (int8)0;
#endif
#endif
    const int inp_offset = mul24(mad24(mad24(mad24(b_idx,in_channel_block,c_idx),input_shape.x,ih_start),src_width,iw_start+input_pad_left),16);
    for(int kh=0; kh<KERNEL_Y; kh++) {
        int ih_cur = ih_start + kh;
        if(ih_cur < 0 || ih_cur >= input_shape.x) {
            continue;
        }

        FLOAT line_cache[INPUT_LINE_SIZE];
        for (int i = 0; i < INPUT_LINE_SIZE; i++) {
            if ((iw_start + i) >= 0 && (iw_start + i) < input_shape.y){
#ifdef MNN_SUPPORT_FP16
                line_cache[i] = as_half(intel_sub_group_block_read_us((__global ushort*)(input + inp_offset + mul24(mad24(kh,src_width,i),16))));
#else
                line_cache[i] = as_float(intel_sub_group_block_read((__global uint*)(input + inp_offset + mul24(mad24(kh,src_width,i),16))));
#endif
            } else{
#ifdef POOL_AVG
                line_cache[i] = 0;
#else
                line_cache[i] = (FLOAT)(-FLT_MAX);
#endif
            }
        }


        for(int kw=0; kw<KERNEL_X; kw++) {
            COMPUTE_FLOAT8 src;
            __attribute__((opencl_unroll_hint(8)))
            for (int i = 0; i < 8; i++) {
                src[i] = line_cache[kw + STRIDE_X*i];
            }
#ifdef POOL_AVG
            result += src;
#else
#if RETURN_REDICE
            redice = src > result ? (int8)((ih_start + kh) * input_shape.y + iw_start + kw) : redice;
#endif
            result = fmax(result, src);
#endif
        }
    }
#ifdef POOL_AVG
    result = result / total_count;
#endif


    const uint lid_x = sglid % 4;
    const uint lid_y = sglid / 4;
    
    const int out_offset = (((b_idx + c_idx * 4 * batch)*output_shape.x + oh_idx)* output_shape.y + ow_idx + output_pad_left)*4;
    const int batch_width_height = batch * output_shape.y * output_shape.x * 4;
#if RETURN_REDICE
    const int redice_offset = (((b_idx + c_idx * 4 * batch)*output_shape.x + oh_idx)* output_shape.y + ow_idx)*4;
#endif
#if OUTPUT_LEFTOVERS
    if ((c_idx+1)*16 >= channel) {
        for (int i = 0; i < 8; i++) {
            if ((c_idx*16 + lid_y * 4 + lid_x < channel) && (ow_idx + i) < output_shape.y)
                output[out_offset + lid_y * batch_width_height + i * 4 + lid_x] = result[i];
#if RETURN_REDICE
                rediceOutput[redice_offset + lid_y * batch_width_height + i * 4 + lid_x] = redice[i];
#endif
        }
    }
    else
#endif  
    {
        for (int i = 0; i < 8 && (ow_idx + i) < output_shape.y; i++) {
            output[out_offset + lid_y * batch_width_height + i * 4 + lid_x] = result[i];
#if RETURN_REDICE
            rediceOutput[redice_offset + lid_y * batch_width_height + i * 4 + lid_x] = redice[i];
#endif
        }
    }
}