
#include "bt.h"

__kernel void rhs_norm(__global double *g_rhs,
                       __global double *g_rms,
                       __local double *l_rms,
                       int gp0,
                       int gp1,
                       int gp2)
{
  int i, j, k, m, lid;
  double add;
  __local double *rms_local;

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

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

  if (k <= (gp2-2)) {
    __global double (*rhs)[JMAXP+1][IMAXP+1][5] = 
      (__global double (*)[JMAXP+1][IMAXP+1][5])g_rhs;

    for (j = 1; j <= (gp1-2); j++) {
      for (i = 1; i <= (gp0-2); i++) {
        for (m = 0; m < 5; m++) {
          add = rhs[k][j][i][m];
          rms_local[m] = rms_local[m] + add*add;
        } 
      } 
    } 
  }

  barrier(CLK_LOCAL_MEM_FENCE);

  if (lid == 0) {
    for (i = 1; i < get_local_size(0); i++) {
      __local double *rms_other = &l_rms[i * 5];
      for (m = 0; m < 5; m++) {
        rms_local[m] += rms_other[m];
      }
    }

    __global double *rms = &g_rms[get_group_id(0) * 5];
    for (m = 0; m < 5; m++) {
      rms[m] = rms_local[m];
    }
  }
}

