#define SUM(a, b) (a + b)
#define MUL(a, b) (a * b)
#define MIN(a, b) min(a, b)
#define MAX(a, b) max(a, b)

__kernel
void greduce(
            __global float* buffer,
            __global float* result,
            __local float* scratch,
            __const int length) {

  int global_index = get_global_id(0);
  int local_index = get_local_id(0);
  
  // Load data into local memory
  if (global_index < length) {
    scratch[local_index] = buffer[global_index];
  } else {
    // Infinity is the identity element for the sum operation
    scratch[local_index] = NEUTRAL;
  }
  
  barrier(CLK_LOCAL_MEM_FENCE);
  
  for(int offset = 1; offset < get_local_size(0); offset <<= 1) {
    int mask = (offset << 1) - 1;
    if ((local_index & mask) == 0) {
      float other = scratch[local_index + offset];
      float mine = scratch[local_index];
      scratch[local_index] = OP(mine, other);
    }
    
    barrier(CLK_LOCAL_MEM_FENCE);
  }
  
  if (local_index == 0) {
    result[get_group_id(0)] = scratch[0];
  }
}