 #include <bits/stdc++.h>
-
+#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
 using namespace std;
 
 typedef long long int LL;

  }
}

vector<LL> fft_con(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);

  LL p_A[A.size()];
  LL p_powers[powers.size()];
  copy(A.begin(), A.end(), p_A);
  copy(powers.begin(), powers.end(), p_powers);

  LL *d_A, *d_powers;
  cudaMalloc(&d_A, A.size() * sizeof(LL));
  cudaMalloc(&d_powers, powers.size() * sizeof(LL));

  cudaMemcpy (d_A, p_A, A.size() * sizeof (LL), cudaMemcpyHostToDevice);
  cudaMemcpy (d_powers, p_powers, powers.size() * sizeof (LL), cudaMemcpyHostToDevice);

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

  fft_kernel<<<dimGrid, dimBlock>>> (d_A, dir, prime, ln, d_powers, A.size());

  cudaDeviceSynchronize();

  A.clear();
  cudaMemcpy (p_A, d_A, a.size() * sizeof (LL), cudaMemcpyDeviceToHost);
  copy(&p_A[0], &p_A[a.size()], back_inserter(A));

  cudaFree(d_A);
  cudaFree(d_powers);


  return A;

}

int main(){
  LL prime = ROU_2[0].first;
  LL basew = ROU_2[0].second;
  vector<LL> a(4096);
  for (int i = 0; i < a.size(); i++){
    a[i] = (i + 2) * 4;
  }

  vector<LL> A = fft(a, 1, prime, basew);
  vector<LL> B = fft_con(a, 1, prime, basew);

  if (A == B)
    cout << "yay!" << endl;

  return 0;
}

            make_pair(754974721,643797295), make_pair(1107296257,883865065)};

PLL ROU_2[] = {make_pair(1711276033LL, 1223522572LL), make_pair(1790967809LL, 1110378081LL)};

inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

PLL ext_euclid(LL a, LL b) {
  if (b == 0)
