__kernel void reduce_min(__global const int* A, __global int* B, __local int* scratch) {
        //Get local variable data
        int id = get_global_id(0);
        int lid = get_local_id(0);
        int N = get_local_size(0);

        //Store valus of global memory into local memory
        scratch[lid] = A[id];

        //Wait for copying to complete
        barrier(CLK_LOCAL_MEM_FENCE);

        for (int i = 1; i < N; i *= 2) {
                if (!(lid % (i * 2)) && ((lid + i) < N)) {
                        if (scratch[lid] > scratch[lid + i])
                                scratch[lid] = scratch[lid+i];
                }

                barrier(CLK_LOCAL_MEM_FENCE);
        }

        //Store cache in output array
        if (!lid)
                atomic_min(&B[0], scratch[lid]);
}


__kernel void reduce_max(__global const int* A, __global int* B, __local int* scratch) {
        //Get local variable data
        int id = get_global_id(0);
        int lid = get_local_id(0);
        int N = get_local_size(0);

        //Store valus of global memory into local memory
        scratch[lid] = A[id];

        //Wait for copying to complete
        barrier(CLK_LOCAL_MEM_FENCE);

        for (int i = 1; i < N; i *= 2) {
                if (!(lid % (i * 2)) && ((lid + i) < N)) {
                        if (scratch[lid] < scratch[lid + i])
                                scratch[lid] = scratch[lid+i];
                }

                barrier(CLK_LOCAL_MEM_FENCE);
        }

        //Store cache in output array
        if (!lid)
                atomic_max(&B[0], scratch[lid]);
}

__kernel void reduce_avg(__global const int* A, __global int* B, __local int* scratch) {
        //Get local variable data
        int id = get_global_id(0);
        int lid = get_local_id(0);
        int N = get_local_size(0);

        //Store valus of global memory into local memory
        scratch[lid] = A[id];

        //Wait for copying to complete
        barrier(CLK_LOCAL_MEM_FENCE);

        for (int i = 1; i < N; i *= 2) {
                if (!(lid % (i * 2)) && ((lid + i) < N))
                {
                        scratch[lid] += scratch[lid+i];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
        }

        //Store cache in output array
        if (!lid)
                atomic_add(&B[0],scratch[lid]);
}