#include "lu.h"

__kernel void pintgr_reduce(__global double *g_phi1,
                            __global double *g_phi2,
                            __global double *g_frc,
                            __local double *l_frc,
                            int ibeg,
                            int ifin1,
                            int jbeg,
                            int jfin1)
{
  int i, j, lid;
  double my_frc = 0.0;

  j = get_global_id(0) + jbeg;
  lid = get_local_id(0);

  if (j < jfin1) {
    __global double (*phi1)[ISIZ2+2] = (__global double (*)[ISIZ2+2])g_phi1;
    __global double (*phi2)[ISIZ2+2] = (__global double (*)[ISIZ2+2])g_phi2;

    for (i = ibeg; i < ifin1; i++) {
      my_frc = my_frc + (  phi1[j][i]
                         + phi1[j][i+1]
                         + phi1[j+1][i]
                         + phi1[j+1][i+1]
                         + phi2[j][i]
                         + phi2[j][i+1]
                         + phi2[j+1][i]
                         + phi2[j+1][i+1] );
    }
  }
  l_frc[lid] = my_frc;

  barrier(CLK_LOCAL_MEM_FENCE);

  if (lid == 0) {
    for (i = 1; i < get_local_size(0); i++) {
      my_frc += l_frc[i];
    }

    g_frc[get_group_id(0)] = my_frc;
  }
}

