#include "lu.h"

__kernel void l2norm(__global double *g_v,
                     __global double *g_sum,
                     __local double *l_sum,
                     int l_nz0,
                     int l_ist,
                     int l_iend,
                     int l_jst,
                     int l_jend)
{
  int i, j, k, m, lid;
  __global double (*v)[ISIZ2/2*2+1][ISIZ1/2*2+1][5];
  __local double *sum_local;

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

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

  if (k < nz0-1) {
    v = (__global double (*)[ISIZ2/2*2+1][ISIZ1/2*2+1][5])g_v;

    for (j = l_jst; j < l_jend; j++) {
      for (i = l_ist; i < l_iend; i++) {
        for (m = 0; m < 5; m++) {
          sum_local[m] = sum_local[m] + v[k][j][i][m] * v[k][j][i][m];
        }
      }
    }
  }
  
  barrier(CLK_LOCAL_MEM_FENCE);

  if (lid == 0) {
    for (i = 1; i < get_local_size(0); i++) {
      __local double *sum_other = &l_sum[i * 5];
      for (m = 0; m < 5; m++) {
        sum_local[m] += sum_other[m];
      }
    }

    __global double *sum = &g_sum[get_group_id(0) * 5];
    for (m = 0; m < 5; m++) {
      sum[m] = sum_local[m];
    }
  }
}

