1#ifndef AMREX_FFT_LOCAL_R2C_H_
2#define AMREX_FFT_LOCAL_R2C_H_
3#include <AMReX_Config.H>
33 int M = AMREX_SPACEDIM>
60 bool cache_plan =
true);
62 bool cache_plan =
false);
135 return m_spectral_size;
143 T* m_p_fwd =
nullptr;
146#if defined(AMREX_USE_SYCL)
153 bool m_cache_plan =
false;
156template <
typename T, FFT::Direction D,
int M>
161 m_real_size(fft_size),
162 m_spectral_size(fft_size)
163#if defined(AMREX_USE_GPU)
164 , m_cache_plan(cache_plan)
167#if !defined(AMREX_USE_GPU)
172 m_spectral_size[0] = m_real_size[0]/2 + 1;
174#if defined(AMREX_USE_SYCL)
182 auto* pf = (
void*)m_p_fwd;
183 auto* pb = (
void*)m_p_bwd;
186 m_fft_fwd.template init_r2c<Direction::forward,M>(m_real_size, pf, pb, m_cache_plan);
187 m_fft_bwd = m_fft_fwd;
190 m_fft_fwd.template init_r2c<Direction::forward,M>(m_real_size, pf, pb, m_cache_plan);
193 m_fft_bwd.template init_r2c<Direction::backward,M>(m_real_size, pf, pb, m_cache_plan);
197#if defined(AMREX_USE_SYCL)
202template <
typename T, FFT::Direction D,
int M>
206 if (m_fft_bwd.plan != m_fft_fwd.plan) {
216template <
typename T, FFT::Direction D,
int M>
219 static_assert(M >= 1 && M <= 3);
223template <
typename T, FFT::Direction D,
int M>
225 : m_p_fwd(rhs.m_p_fwd),
226 m_p_bwd(rhs.m_p_bwd),
227 m_fft_fwd(rhs.m_fft_fwd),
228 m_fft_bwd(rhs.m_fft_bwd),
229#if defined(AMREX_USE_SYCL)
230 m_gpu_stream(rhs.m_gpu_stream),
232 m_real_size(rhs.m_real_size),
233 m_spectral_size(rhs.m_spectral_size),
234 m_cache_plan(rhs.m_cache_plan)
236 rhs.m_cache_plan =
true;
239template <
typename T, FFT::Direction D,
int M>
242 if (
this == &rhs) {
return *
this; }
246 m_p_fwd = rhs.m_p_fwd;
247 m_p_bwd = rhs.m_p_bwd;
248 m_fft_fwd = rhs.m_fft_fwd;
249 m_fft_bwd = rhs.m_fft_bwd;
250#if defined(AMREX_USE_SYCL)
251 m_gpu_stream = rhs.m_gpu_stream;
253 m_real_size = rhs.m_real_size;
254 m_spectral_size = rhs.m_spectral_size;
255 m_cache_plan = rhs.m_cache_plan;
257 rhs.m_cache_plan =
true;
262template <
typename T, FFT::Direction D,
int M>
269#if defined(AMREX_USE_GPU)
271 m_fft_fwd.set_ptrs((
void*)indata, (
void*)outdata);
273#if defined(AMREX_USE_SYCL)
275 if (current_stream != m_gpu_stream) {
283 if (((T*)indata != m_p_fwd) || (outdata != m_p_bwd)) {
284 m_p_fwd = (T*)indata;
286 auto* pf = (
void*)m_p_fwd;
287 auto* pb = (
void*)m_p_bwd;
289 m_fft_fwd.template init_r2c<Direction::forward,M>(m_real_size, pf, pb,
false);
292 m_fft_bwd.template init_r2c<Direction::backward,M>(m_real_size, pf, pb,
false);
298 m_fft_fwd.template compute_r2c<Direction::forward>();
300#if defined(AMREX_USE_SYCL)
301 if (current_stream != m_gpu_stream) {
307template <
typename T, FFT::Direction D,
int M>
314#if defined(AMREX_USE_GPU)
316 m_fft_bwd.set_ptrs((
void*)outdata, (
void*)indata);
318#if defined(AMREX_USE_SYCL)
320 if (current_stream != m_gpu_stream) {
328 if (((
GpuComplex<T>*)indata != m_p_bwd) || (outdata != m_p_fwd)) {
331 auto* pf = (
void*)m_p_fwd;
332 auto* pb = (
void*)m_p_bwd;
334 m_fft_bwd.template init_r2c<Direction::backward,M>(m_real_size, pf, pb,
false);
337 m_fft_fwd.template init_r2c<Direction::forward,M>(m_real_size, pf, pb,
false);
343 m_fft_bwd.template compute_r2c<Direction::backward>();
345#if defined(AMREX_USE_SYCL)
346 if (current_stream != m_gpu_stream) {
352template <
typename T, FFT::Direction D,
int M>
356 for (
auto s : m_real_size) {
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
Local Discrete Fourier Transform.
Definition AMReX_FFT_LocalR2C.H:35
void backward(GpuComplex< T > const *indata, T *outdata)
Backward transform.
Definition AMReX_FFT_LocalR2C.H:310
T scalingFactor() const
Scaling factor for normalization.
Definition AMReX_FFT_LocalR2C.H:353
void forward(T const *indata, GpuComplex< T > *outdata)
Forward transform.
Definition AMReX_FFT_LocalR2C.H:265
void clear()
Destroy any cached FFT plans; leaves the object in an uninitialized state.
Definition AMReX_FFT_LocalR2C.H:203
LocalR2C & operator=(LocalR2C &&) noexcept
Move assignment; releases current plans and adopts those from rhs.
Definition AMReX_FFT_LocalR2C.H:240
~LocalR2C()
Definition AMReX_FFT_LocalR2C.H:217
IntVectND< M > const & spectralSize() const
Spectral domain extents associated with this plan.
Definition AMReX_FFT_LocalR2C.H:134
static gpuStream_t setStream(gpuStream_t s) noexcept
Definition AMReX_GpuDevice.cpp:735
static void resetStreamIndex() noexcept
Definition AMReX_GpuDevice.H:96
An Integer Vector in dim-Dimensional Space.
Definition AMReX_IntVect.H:57
Definition AMReX_FFT_Helper.H:52
Direction
Definition AMReX_FFT_Helper.H:54
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:263
gpuStream_t gpuStream() noexcept
Definition AMReX_GpuDevice.H:244
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:139
cudaStream_t gpuStream_t
Definition AMReX_GpuControl.H:83
Definition AMReX_FFT_Helper.H:180
A host / device complex number type, because std::complex doesn't work in device code with Cuda yet.
Definition AMReX_GpuComplex.H:30