
#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 cffts1(__global dcomplex *x,
                     __global dcomplex *xout,
                     __global dcomplex *u,
                     int is,
                     int d1,
                     int d2,
                     int d3,
                     int logd1)
{
  __local dcomplex ty1[MAXDIM];
  __local dcomplex ty2[MAXDIM];

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

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

  cfftz(is, logd1, d1, u, ty1, ty2);

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

  for (int i = lid; i < d1; i += LSIZE) {
    int xout_idx = kj_idx + i;
    xout[xout_idx].real = ty1[i].real;
    xout[xout_idx].imag = ty1[i].imag;
  }
}

