link cufft

Grid/algorithms/FFT.h

@@ -36,6 +36,7 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
 #include <hipfft/hipfft.h>
 #endif
 
+#if !defined(GRID_CUDA) && !defined(GRID_HIP)
 #ifdef HAVE_FFTW
 #if defined(USE_MKL) || defined(GRID_SYCL)
 #include <fftw/fftw3.h>
@@ -43,6 +44,7 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
 #include <fftw3.h>
 #endif
 #endif
+#endif
 
 NAMESPACE_BEGIN(Grid);
 
@@ -138,7 +140,7 @@ public:
     
   inline static void fftw_execute_dft(const FFTW_plan p,FFTW_scalar *in,FFTW_scalar *out, int sign) {
     if ( sign == forward ) cufftExecZ2Z(p,in,out,CUFFT_FORWARD);
-    else                   cufftExecZ2Z(p,in,out,CUFFT_BACKWARD);
+    else                   cufftExecZ2Z(p,in,out,CUFFT_INVERSE);
     accelerator_barrier();
   }
   inline static void fftw_destroy_plan(const FFTW_plan p) {
@@ -165,7 +167,7 @@ public:
     
   inline static void fftw_execute_dft(const FFTW_plan p,FFTW_scalar *in,FFTW_scalar *out, int sign) {
     if ( sign == forward ) cufftExecC2C(p,in,out,CUFFT_FORWARD);
-    else                   cufftExecC2C(p,in,out,CUFFT_BACKWARD);
+    else                   cufftExecC2C(p,in,out,CUFFT_INVERSE);
     accelerator_barrier();
   }
   inline static void fftw_destroy_plan(const FFTW_plan p) {
@@ -174,6 +176,7 @@ public:
 };
 #endif
 
+#if !defined(GRID_CUDA) && !defined(GRID_HIP)
 #ifdef HAVE_FFTW
 template<> struct FFTW<ComplexD> {
 public:
@@ -216,6 +219,7 @@ public:
   }
 };
 #endif
+#endif
 
 class FFT {
 private:

Grid/communicator/Communicator_mpi3.cc

@@ -349,7 +349,6 @@ void CartesianCommunicator::SendToRecvFromBegin(std::vector<MpiCommsRequest_t> &
 
   GRID_ASSERT(dest != _processor);
   GRID_ASSERT(from != _processor);
-  GRID_ASSERT(bytes/(sizeof(int32_t))<= 2*1024*1024*1024);
   int tag;
 
   tag= dir+from*32;

configure.ac

@@ -292,13 +292,14 @@ AC_ARG_ENABLE([accelerator],
 case ${ac_ACCELERATOR} in
     cuda)
       echo CUDA acceleration
-      LIBS="${LIBS} -lcuda -lcublas"
+      LIBS="${LIBS} -lcuda -lcublas -lcufft"
       AC_DEFINE([GRID_CUDA],[1],[Use CUDA offload]);;
     sycl)
       echo SYCL acceleration
       AC_DEFINE([GRID_SYCL],[1],[Use SYCL offload]);;
     hip)
       echo HIP acceleration
+      LIBS="${LIBS} -lhipblas -lrocblas -lhipfft"
       AC_DEFINE([GRID_HIP],[1],[Use HIP offload]);;
     none)
       echo NO acceleration    ;;