aboutsummaryrefslogtreecommitdiff
path: root/Source
diff options
context:
space:
mode:
Diffstat (limited to 'Source')
-rw-r--r--Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp23
-rw-r--r--Source/Particles/Collision/CollisionType.cpp4
-rw-r--r--Source/Particles/Deposition/CurrentDeposition.H17
3 files changed, 29 insertions, 15 deletions
diff --git a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp
index dc7f58f48..c98c79835 100644
--- a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp
+++ b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp
@@ -77,8 +77,13 @@ SpectralFieldDataRZ::SpectralFieldDataRZ (amrex::BoxArray const & realspace_ba,
int ostride = grid_size[0];
int odist = 1;
int batch = grid_size[0]; // number of ffts
+# ifdef AMREX_USE_FLOAT
+ auto cufft_type = CUFFT_C2C;
+# else
+ auto cufft_type = CUFFT_Z2Z;
+# endif
result = cufftPlanMany(&forward_plan[mfi], 1, fft_length, inembed, istride, idist,
- onembed, ostride, odist, CUFFT_Z2Z, batch);
+ onembed, ostride, odist, cufft_type, batch);
if (result != CUFFT_SUCCESS) {
amrex::AllPrint() << " cufftPlanMany failed! \n";
}
@@ -225,9 +230,13 @@ SpectralFieldDataRZ::FABZForwardTransform (amrex::MFIter const & mfi, amrex::Box
cudaStream_t stream = amrex::Gpu::Device::cudaStream();
cufftSetStream(forward_plan[mfi], stream);
for (int mode=0 ; mode < n_rz_azimuthal_modes ; mode++) {
+# ifdef AMREX_USE_FLOAT
+ result = cufftExecC2C(forward_plan[mfi],
+# else
result = cufftExecZ2Z(forward_plan[mfi],
- reinterpret_cast<cuDoubleComplex*>(tempHTransformed[mfi].dataPtr(mode)), // cuDoubleComplex *in
- reinterpret_cast<cuDoubleComplex*>(tmpSpectralField[mfi].dataPtr(mode)), // cuDoubleComplex *out
+# endif
+ reinterpret_cast<AnyFFT::Complex*>(tempHTransformed[mfi].dataPtr(mode)), // Complex *in
+ reinterpret_cast<AnyFFT::Complex*>(tmpSpectralField[mfi].dataPtr(mode)), // Complex *out
CUFFT_FORWARD);
if (result != CUFFT_SUCCESS) {
amrex::AllPrint() << " forward transform using cufftExecZ2Z failed ! \n";
@@ -330,9 +339,13 @@ SpectralFieldDataRZ::FABZBackwardTransform (amrex::MFIter const & mfi, amrex::Bo
cudaStream_t stream = amrex::Gpu::Device::cudaStream();
cufftSetStream(forward_plan[mfi], stream);
for (int mode=0 ; mode < n_rz_azimuthal_modes ; mode++) {
+# ifdef AMREX_USE_FLOAT
+ result = cufftExecC2C(forward_plan[mfi],
+# else
result = cufftExecZ2Z(forward_plan[mfi],
- reinterpret_cast<cuDoubleComplex*>(tmpSpectralField[mfi].dataPtr(mode)), // cuDoubleComplex *in
- reinterpret_cast<cuDoubleComplex*>(tempHTransformed[mfi].dataPtr(mode)), // cuDoubleComplex *out
+# endif
+ reinterpret_cast<AnyFFT::Complex*>(tmpSpectralField[mfi].dataPtr(mode)), // Complex *in
+ reinterpret_cast<AnyFFT::Complex*>(tempHTransformed[mfi].dataPtr(mode)), // Complex *out
CUFFT_INVERSE);
if (result != CUFFT_SUCCESS) {
amrex::AllPrint() << " backwardtransform using cufftExecZ2Z failed ! \n";
diff --git a/Source/Particles/Collision/CollisionType.cpp b/Source/Particles/Collision/CollisionType.cpp
index 36ce7952d..38825d003 100644
--- a/Source/Particles/Collision/CollisionType.cpp
+++ b/Source/Particles/Collision/CollisionType.cpp
@@ -132,7 +132,7 @@ void CollisionType::doCoulombCollisionsWithinTile
#if defined WARPX_DIM_RZ
int ri = (i_cell - i_cell%nz) / nz;
- auto dV = MathConst::pi*(2.0*ri+1.0)*dr*dr*dz;
+ auto dV = MathConst::pi*(2.0_rt*ri+1.0_rt)*dr*dr*dz;
#else
amrex::ignore_unused(nz);
#endif
@@ -229,7 +229,7 @@ void CollisionType::doCoulombCollisionsWithinTile
#if defined WARPX_DIM_RZ
int ri = (i_cell - i_cell%nz) / nz;
- auto dV = MathConst::pi*(2.0*ri+1.0)*dr*dr*dz;
+ auto dV = MathConst::pi*(2.0_rt*ri+1.0_rt)*dr*dr*dz;
#else
amrex::ignore_unused(nz);
#endif
diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H
index 67c6b4bd0..b116bbb64 100644
--- a/Source/Particles/Deposition/CurrentDeposition.H
+++ b/Source/Particles/Deposition/CurrentDeposition.H
@@ -249,12 +249,12 @@ void doDepositionShapeN(const GetParticlePosition& GetPosition,
Complex xy = xy0; // Note that xy is equal to e^{i m theta}
for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
// The factor 2 on the weighting comes from the normalization of the modes
- amrex::Gpu::Atomic::Add( &jx_arr(lo.x+j_jx+ix, lo.y+l_jx+iz, 0, 2*imode-1), 2.*sx_jx[ix]*sz_jx[iz]*wqx*xy.real());
- amrex::Gpu::Atomic::Add( &jx_arr(lo.x+j_jx+ix, lo.y+l_jx+iz, 0, 2*imode ), 2.*sx_jx[ix]*sz_jx[iz]*wqx*xy.imag());
- amrex::Gpu::Atomic::Add( &jy_arr(lo.x+j_jy+ix, lo.y+l_jy+iz, 0, 2*imode-1), 2.*sx_jy[ix]*sz_jy[iz]*wqy*xy.real());
- amrex::Gpu::Atomic::Add( &jy_arr(lo.x+j_jy+ix, lo.y+l_jy+iz, 0, 2*imode ), 2.*sx_jy[ix]*sz_jy[iz]*wqy*xy.imag());
- amrex::Gpu::Atomic::Add( &jz_arr(lo.x+j_jz+ix, lo.y+l_jz+iz, 0, 2*imode-1), 2.*sx_jz[ix]*sz_jz[iz]*wqz*xy.real());
- amrex::Gpu::Atomic::Add( &jz_arr(lo.x+j_jz+ix, lo.y+l_jz+iz, 0, 2*imode ), 2.*sx_jz[ix]*sz_jz[iz]*wqz*xy.imag());
+ amrex::Gpu::Atomic::Add( &jx_arr(lo.x+j_jx+ix, lo.y+l_jx+iz, 0, 2*imode-1), 2._rt*sx_jx[ix]*sz_jx[iz]*wqx*xy.real());
+ amrex::Gpu::Atomic::Add( &jx_arr(lo.x+j_jx+ix, lo.y+l_jx+iz, 0, 2*imode ), 2._rt*sx_jx[ix]*sz_jx[iz]*wqx*xy.imag());
+ amrex::Gpu::Atomic::Add( &jy_arr(lo.x+j_jy+ix, lo.y+l_jy+iz, 0, 2*imode-1), 2._rt*sx_jy[ix]*sz_jy[iz]*wqy*xy.real());
+ amrex::Gpu::Atomic::Add( &jy_arr(lo.x+j_jy+ix, lo.y+l_jy+iz, 0, 2*imode ), 2._rt*sx_jy[ix]*sz_jy[iz]*wqy*xy.imag());
+ amrex::Gpu::Atomic::Add( &jz_arr(lo.x+j_jz+ix, lo.y+l_jz+iz, 0, 2*imode-1), 2._rt*sx_jz[ix]*sz_jz[iz]*wqz*xy.real());
+ amrex::Gpu::Atomic::Add( &jz_arr(lo.x+j_jz+ix, lo.y+l_jz+iz, 0, 2*imode ), 2._rt*sx_jz[ix]*sz_jz[iz]*wqz*xy.imag());
xy = xy*xy0;
}
#endif
@@ -552,8 +552,9 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition,
for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
// The factor 2 comes from the normalization of the modes
// The minus sign comes from the different convention with respect to Davidson et al.
- const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xmin*dxi)*wq*invdtdx/(amrex::Real)imode*
- (sx_new[i]*sz_new[k]*(xy_new - xy_mid) + sx_old[i]*sz_old[k]*(xy_mid - xy_old));
+ const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xmin*dxi)*wq*invdtdx/(amrex::Real)imode
+ *(Complex(sx_new[i]*sz_new[k], 0.)*(xy_new - xy_mid)
+ + Complex(sx_old[i]*sz_old[k], 0.)*(xy_mid - xy_old));
amrex::Gpu::Atomic::Add( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 2*imode-1), djt_cmplx.real());
amrex::Gpu::Atomic::Add( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 2*imode), djt_cmplx.imag());
xy_new = xy_new*xy_new0;
generated by cgit v1.2.3 (git 2.39.1) at 2025-08-08 13:18:29 +0000