
#include "lu.h"

void exact(int i, int j, int k, double u000ijk[], __global double *g_ce);

__kernel void error(__global double *g_u,
                    __global double *g_ce,
                    __global double *g_errnm,
                    __local double *l_errnm,
                    int nx,
                    int ny,
                    int nz)
{
  int i, j, k, m, lid;
  double tmp;
  double u000ijk[5];
  __local double *errnm_local;

  k = get_global_id(0) + 1;
  lid = get_local_id(0);
  errnm_local = &l_errnm[lid * 5];

  for (m = 0; m < 5; m++) {
    errnm_local[m] = 0.0;
  }

  if (k < nz-1) {
    __global double (*u)[ISIZ2/2*2+1][ISIZ1/2*2+1][5] = 
      (__global double (*)[ISIZ2/2*2+1][ISIZ1/2*2+1][5])g_u;

    for (j = jst; j < jend; j++) {
      for (i = ist; i < iend; i++) {
        exact( i, j, k, u000ijk, g_ce );
        for (m = 0; m < 5; m++) {
          tmp = ( u000ijk[m] - u[k][j][i][m] );
          errnm_local[m] = errnm_local[m] + tmp * tmp;
        }
      }
    }
  }

  barrier(CLK_LOCAL_MEM_FENCE);

  if (lid == 0) {
    for (i = 1; i < get_local_size(0); i++) {
      __local double *errnm_other = &l_errnm[i * 5];
      for (m = 0; m < 5; m++) {
        errnm_local[m] += errnm_other[m];
      }
    }

    __global double *errnm = &g_errnm[get_group_id(0) * 5];
    for (m = 0; m < 5; m++) {
      errnm[m] = errnm_local[m];
    }
  }
}

