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

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

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

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

  cfftz(is, logd3, d3, u, ty1, ty2);

  CFFTS_BARRIER(CLK_LOCAL_MEM_FENCE);

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


