
#include "ft.h"

#define CFFTS_BARRIER(m)    barrier(m)

void cfftz(int is, int m, int n, 
           __global dcomplex *u, __local dcomplex *x, __local dcomplex *y);


__kernel void cffts2(__global dcomplex *x,
                     __global dcomplex *xout,
                     __global dcomplex *u,
                     int is,
                     int d1,
                     int d2,
                     int d3,
                     int logd2)
{
  __local dcomplex ty1[MAXDIM];
  __local dcomplex ty2[MAXDIM];

  int k = get_global_id(1);
  int i = get_group_id(0);
  int lid = get_local_id(0);

  int ki_idx = k*d2*(d1+1) + i;
  for (int j = lid; j < d2; j += LSIZE) {
    int x_idx = ki_idx + j*(d1+1);
    ty1[j].real = x[x_idx].real;
    ty1[j].imag = x[x_idx].imag;
  }

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

  cfftz(is, logd2, d2, u, ty1, ty2);

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

  for (int j = lid; j < d2; j += LSIZE) {
    int xout_idx = ki_idx + j*(d1+1);
    xout[xout_idx].real = ty1[j].real;
    xout[xout_idx].imag = ty1[j].imag;
  }
}
