   return ans;
 }
 
-void bit_reverse_copy(vector<LL> &a, vector<LL> &A, int n) {
-  A.resize(a.size());
-  for (int i = 0; i < a.size(); i++)
+void bit_reverse_copy(LL *a, LL *A, int n, int size) {
+  for (int i = 0; i < size; i++)
     A[bit_reverse(i, n)] = a[i];
 }
 
-vector<LL> compute_powers(int ln, LL basew, LL prime){
-  vector<LL> powers(ln);
+void compute_powers(LL *powers, int ln, LL basew, LL prime){
   powers[0] = basew;
   for (int i = 1; i < ln; i++){
     powers[i] = (powers[i - 1] * powers[i - 1]) % prime;
   }
-  return powers;
-}
-
-vector<LL> fft(vector<LL> &a, int dir, LL prime, LL basew) {
-  int ln = ceil(log2(float(a.size())));
-  vector<LL> A;
-  bit_reverse_copy(a, A, ln);
-  vector<LL> powers = compute_powers(ln, basew, prime);
+}
+
+void fft(LL *a, LL *A, int dir, LL prime, LL basew, int size) {
+  int ln = ceil(log2(float(size)));
+  bit_reverse_copy(a, A, ln, size);
+  LL *powers = (LL*) malloc (sizeof (LL) * ln);
+  compute_powers(powers, ln, basew, prime);
 
   for (int s = 1; s <= ln; s++) {
     long long m = (1LL << s);
     LL wm = powers[ln - s];
-    //cout << wm << endl;
     if (dir == -1)
       wm =  mod_inv(wm, prime);
 
-    for (int k = 0; k < a.size(); k += m) {
+    for (int k = 0; k < size; k += m) {
       LL w = 1, mh = m >> 1;
       for (int j = 0; j < mh; j++) {
         LL t = (w * A[k + j + mh]) % prime;

  }

  if (dir < 0) {
    LL in = mod_inv(A.size(), prime);
    for (int i = 0; i < A.size(); i++)
      A[i] = (A[i] * in) % prime;
  }

  return A;
}

__device__ void d_ext_euclid(LL a, LL b, LL &x, LL &y, LL &g) {

  }
}

void fft_con(LL *a, LL *A, int dir, LL prime, LL basew, int size){
  int ln = ceil(log2(float(size)));
  bit_reverse_copy(a, A, ln, size);
  LL *powers = (LL*) malloc (sizeof (LL) * ln);
  compute_powers(powers, ln, basew, prime);

  LL *d_A, *d_powers;
  cudaMalloc(&d_A, size * sizeof(LL));
  cudaMalloc(&d_powers, ln * sizeof(LL));

  cudaMemcpy (d_A, A, size * sizeof (LL), cudaMemcpyHostToDevice);
  cudaMemcpy (d_powers, powers, ln * sizeof (LL), cudaMemcpyHostToDevice);

  dim3 dimGrid(ceil(float(size / 1024.0)), 1, 1);
  dim3 dimBlock(1024, 1, 1);

  fft_kernel<<<dimGrid, dimBlock>>> (d_A, dir, prime, ln, d_powers, size);
  gpuErrchk( cudaPeekAtLastError() );
  gpuErrchk( cudaDeviceSynchronize() );

  cudaDeviceSynchronize();

  cudaMemcpy (A, d_A, size * sizeof (LL), cudaMemcpyDeviceToHost);

  cudaFree(d_A);
  cudaFree(d_powers);
  free(powers);

}

bool cmp_vectors (LL *A, LL *B, int size){
  for (int i = 0; i < size; i++){
    if (A[i] != B[i]){
      cout << A[i] << " " << B[i] << " i: " << i << endl;
      return false;
    }
  }
  return true;
}

int main(){
  LL prime = ROU_2[0].first;
  LL basew = ROU_2[0].second;
  int size = 8192;
  LL *a = (LL*) malloc( sizeof (LL) * size);
  LL *A = (LL*) malloc( sizeof (LL) * size);
  LL *B = (LL*) malloc( sizeof (LL) * size);
  for (int i = 0; i < size; i++){
    a[i] = i;
  }

  fft(a, A, 1, prime, basew, size);
  fft_con(a, B, 1, prime, basew, size);

  if (!cmp_vectors (A, B, size))
    cout << "nay :(" << endl;
  else
    cout << "yay :)" << endl;

  free(a);
  free(A);
  free(B);
  return 0;
}
