__kernel void pixelArtUpscale(__global const int* inputImage, __global int* outputImage, const int sourceWidth){

    int lx = get_local_id(0) + 1;
    int ly = get_local_id(1) + 1;

    int gx = get_global_id(0);
    int gy = get_global_id(1);

    __local unsigned int localBuffer[32 + 2][32 + 2];

    const size_t rowcol = (get_global_id(1) * get_global_size(0) + get_global_id(0));
    localBuffer[lx][ly] = inputImage[rowcol];

    // load pixel left and right when at the edge
    if (lx == 1){
        if (rowcol % get_global_size(0) == 0){
            localBuffer[lx - 1][ly] = inputImage[rowcol];
        } else{
            localBuffer[lx - 1][ly] = inputImage[rowcol - 1];
        }
    }
    else if (lx == get_local_size(0)){
        if (rowcol % get_global_size(0) == get_global_size(0) - 1){
            localBuffer[lx + 1][ly] = inputImage[rowcol];
        } else{
            localBuffer[lx + 1][ly] = inputImage[rowcol + 1];
        }
    }

    // load pixel top and bottom when at the edge
    if (ly == 1){
        if (rowcol >= get_global_size(0)) {
            localBuffer[lx][ly - 1] = inputImage[rowcol - get_global_size(0)];
        } else {
            localBuffer[lx][ly - 1] = inputImage[rowcol];
        }
    }
    else if (ly == get_local_size(1)){
        if (rowcol < get_global_size(0)*get_global_size(1) - get_global_size(0)){
            localBuffer[lx][ly + 1] = inputImage[rowcol + get_global_size(0)];
        } else {
            localBuffer[lx][ly + 1] = inputImage[rowcol];
        }
    }

    // load 4 corner pixel
    const size_t Localrowcol = (get_local_id(1) * get_local_size(0) + get_local_id(0));
    if (Localrowcol == 0) {
        localBuffer[lx - 1][ly - 1] = inputImage[rowcol];
    }
    if (Localrowcol == get_local_size(0) - 1) {
        localBuffer[lx + 1][ly - 1] = inputImage[rowcol];
    }
    if (Localrowcol == get_local_size(0)*get_local_size(1) - get_local_size(0)){
        localBuffer[lx - 1][ly + 1] = inputImage[rowcol];
    }
    if (Localrowcol == get_local_size(0)*get_local_size(1) - 1){
        localBuffer[lx + 1][ly + 1] = inputImage[rowcol];
    }

    // wait on all to be filled
    barrier(CLK_LOCAL_MEM_FENCE);


    int A = localBuffer[lx - 1][ly - 1];
    int B = localBuffer[lx][ly - 1];
    int C = localBuffer[lx + 1][ly - 1];
    int D = localBuffer[lx - 1][ly - 1];
    int E = localBuffer[lx][ly];
    int F = localBuffer[lx + 1][ly];
    int G = localBuffer[lx - 1][ly + 1];
    int H = localBuffer[lx][ly + 1];
    int I = localBuffer[lx + 1][ly + 1];

    //Eagle impleentation
    int E00 = E;
    int E10 = E;
    int E01 = E;
    int E11 = E;

    if(D == A && A == B) E00 = A;
    if(B == C && C == F) E01 = C;
    if(D == G && G == H) E10 = G;
    if(F == I && I == H) E11 = I;


    outputImage[2 * gy * (2 * get_global_size(0)) + 2 * gx] = E00;
    outputImage[ 2* gy * (2 * get_global_size(0)) + 2 * gx + 1] = E01;
    outputImage[(2 * gy + 1) * (2 * get_global_size(0)) + 2 * gx] = E10;
    outputImage[(2 * gy + 1) * (2 * get_global_size(0)) + 2 * gx + 1] = E11;

}