MinMax-Kernel

#define WORKGROUPSIZE 64
kernel void MinMax( global float * input,
            global float * output,
                        int sizex, //size in x
            int sizey, //size in y
            int sizez, //size in z
            int size //global size x*y*z
                    )
{
    //get global ids
        int idx = get_global_id(0);
        int idy = get_global_id(1);
        int idz = get_global_id(2);

    //get local ids
        int idx_local = get_local_id(0);
        int idy_local = get_local_id(1);
        int idz_local = get_local_id(2);
    //get sizes of workgroup
        int local_x = get_local_size(0);
        int local_y = get_local_size(1);
        int local_z = get_local_size(2);
    //calculate size of workgroup
        int local_size = local_x*local_y*local_z;
    //because array is saved as 1D, recalculate index
        int gid = idx + sizex*idy + sizex*sizey*idz;
    //same for work-group
        int gid_local = idx_local + local_x*idy_local +local_x*local_y*idz_local;
    //get work-group count
        int group_count = get_num_groups(0) * get_num_groups(1) * get_num_groups(2);

    //local memory init
    //maybe I should put that in the arguments
        __local float Minvals[ WORKGROUPSIZE ];
        __local float Maxvals[ WORKGROUPSIZE ];

    //make dirty memory undirty^^
    if(gid_local==0)
    {
        for(int i =0; i<WORKGROUPSIZE;i++)
        {
            Minvals[i] = MAXFLOAT;
            Maxvals[i] = 0;
        }
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    //copy from global to local
    if(gid < size)
    {
        Minvals[gid_local] = input[gid];
        Maxvals[gid_local] = input[gid];
    }

    barrier(CLK_LOCAL_MEM_FENCE);
        //local parallel reduction
    for(int offset = local_size/2;offset>0;offset/=2)
    {
        barrier(CLK_LOCAL_MEM_FENCE);
        if(gid_local<offset)
        {
            Minvals[gid_local] = fmin(Minvals[gid_local],Minvals[gid_local+offset]);
            Maxvals[gid_local] = fmax(Maxvals[gid_local],Maxvals[gid_local+offset]);
        }

    }

    barrier(CLK_GLOBAL_MEM_FENCE);
    //Global write back
    if(gid_local == 0)
    {
        output[group_id*2] =  Minvals[gid_local];
        output[group_id*2 +1] = Maxvals[gid_local];
    }

    barrier(CLK_GLOBAL_MEM_FENCE);
    //global reduction
    if(gid==0)
    {
        float mint=256;
        float maxt=0;
        for(int i =0; i < group_count; i=i+2)
        {
            mint     = fmin(mint,output[i]);
            maxt     = fmax(maxt,output[i+1]);

        }
        //write to output -> end
        output[0] = mint;
        output[1] = maxt;
    }

}