#include "ft.h"

#define CFFTS_BARRIER(m)    barrier(m)


__kernel void checksum(__global dcomplex *u1,
                       __global dcomplex *g_chk,
                       __local dcomplex *l_chk,
                       int d1,
                       int d2)
{
  int q, r, s;
  int j = get_global_id(0) + 1;
  int lid = get_local_id(0);

  if (j <= 1024) {
    q = j % NX;
    r = 3*j % NY;
    s = 5*j % NZ;
    int u1_idx = s*d2*(d1+1) + r*(d1+1) + q;
    l_chk[lid].real = u1[u1_idx].real;
    l_chk[lid].imag = u1[u1_idx].imag;
  } else {
    l_chk[lid] = dcmplx(0.0, 0.0);
  }

  barrier(CLK_LOCAL_MEM_FENCE);

  // reduction in a work-group
  for (int i = get_local_size(0) / 2; i > 0; i >>= 1) {
    if (lid < i) {
      l_chk[lid] = dcmplx_add(l_chk[lid], l_chk[lid+i]);
    }
    barrier(CLK_LOCAL_MEM_FENCE);
  }
  if (lid == 0) {
    g_chk[get_group_id(0)].real = l_chk[0].real;
    g_chk[get_group_id(0)].imag = l_chk[0].imag;
  }
}

