diff options
author | 2018-10-26 10:48:23 -0700 | |
---|---|---|
committer | 2018-10-26 10:48:23 -0700 | |
commit | f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c (patch) | |
tree | 78f1cc8ae3ed551d4b3dd76b72e589d42561eb46 | |
parent | 8642045c0f27a087430a0a1083db7fb678e8fe2a (diff) | |
download | WarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.tar.gz WarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.tar.zst WarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.zip |
update according to amrex changes
-rw-r--r-- | Source/PhysicalParticleContainer.H | 2 | ||||
-rw-r--r-- | Source/PhysicalParticleContainer.cpp | 26 | ||||
-rw-r--r-- | Source/WarpXEvolve.cpp | 12 | ||||
-rw-r--r-- | Source/WarpXMove.cpp | 4 |
4 files changed, 22 insertions, 22 deletions
diff --git a/Source/PhysicalParticleContainer.H b/Source/PhysicalParticleContainer.H index 16c1c6141..42ccdf065 100644 --- a/Source/PhysicalParticleContainer.H +++ b/Source/PhysicalParticleContainer.H @@ -111,7 +111,7 @@ protected: const amrex::RealBox& particle_real_box); void AddPlasmaCPU (int lev, amrex::RealBox part_realbox); -#ifdef AMREX_USE_CUDA +#ifdef AMREX_USE_GPU void AddPlasmaGPU (int lev, amrex::RealBox part_realbox); #endif int GetRefineFac(const amrex::Real x, const amrex::Real y, const amrex::Real z); diff --git a/Source/PhysicalParticleContainer.cpp b/Source/PhysicalParticleContainer.cpp index d74961437..0ba081278 100644 --- a/Source/PhysicalParticleContainer.cpp +++ b/Source/PhysicalParticleContainer.cpp @@ -257,7 +257,7 @@ PhysicalParticleContainer::AddParticles (int lev) void PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) { -#ifdef AMREX_USE_CUDA +#ifdef AMREX_USE_GPU AddPlasmaGPU(lev, part_realbox); #else AddPlasmaCPU(lev, part_realbox); @@ -480,7 +480,7 @@ PhysicalParticleContainer::AddPlasmaCPU (int lev, RealBox part_realbox) } } -#ifdef AMREX_USE_CUDA +#ifdef AMREX_USE_GPU void PhysicalParticleContainer::AddPlasmaGPU (int lev, RealBox part_realbox) { @@ -1334,7 +1334,7 @@ PhysicalParticleContainer::Evolve (int lev, const int ncomp = 1; FArrayBox const* local_fab = local_rho[thread_num].get(); FArrayBox* global_fab = &rhofab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tile_box, tbx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tile_box, tbx, { global_fab->atomicAdd(*local_fab, tbx, tbx, 0, icomp, ncomp); }); @@ -1384,7 +1384,7 @@ PhysicalParticleContainer::Evolve (int lev, const int ncomp = 1; FArrayBox const* local_fab = local_rho[thread_num].get(); FArrayBox* global_fab = &crhofab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tile_box, tbx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tile_box, tbx, { global_fab->atomicAdd(*local_fab, tbx, tbx, 0, icomp, ncomp); }); @@ -1562,19 +1562,19 @@ PhysicalParticleContainer::Evolve (int lev, jz_ptr = local_jz[thread_num]->dataPtr(); FArrayBox* local_jx_ptr = local_jx[thread_num].get(); - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, b, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbx, b, { local_jx_ptr->setVal(0.0, b, 0, 1); }); FArrayBox* local_jy_ptr = local_jy[thread_num].get(); - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tby, b, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tby, b, { local_jy_ptr->setVal(0.0, b, 0, 1); }); FArrayBox* local_jz_ptr = local_jz[thread_num].get(); - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbz, b, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbz, b, { local_jz_ptr->setVal(0.0, b, 0, 1); }); @@ -1605,21 +1605,21 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox const* local_jx_const_ptr = local_jx[thread_num].get(); FArrayBox* global_jx_ptr = &jxfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, { global_jx_ptr->atomicAdd(*local_jx_const_ptr, thread_bx, thread_bx, 0, 0, 1); }); FArrayBox const* local_jy_const_ptr = local_jy[thread_num].get(); FArrayBox* global_jy_ptr = &jyfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tby, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tby, thread_bx, { global_jy_ptr->atomicAdd(*local_jy_const_ptr, thread_bx, thread_bx, 0, 0, 1); }); FArrayBox const* local_jz_const_ptr = local_jz[thread_num].get(); FArrayBox* global_jz_ptr = &jzfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbz, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbz, thread_bx, { global_jz_ptr->atomicAdd(*local_jz_const_ptr, thread_bx, thread_bx, 0, 0, 1); }); @@ -1683,21 +1683,21 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox const* local_jx_ptr = local_jx[thread_num].get(); FArrayBox* global_jx_ptr = &cjxfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, { global_jx_ptr->atomicAdd(*local_jx_ptr, thread_bx, thread_bx, 0, 0, 1); }); FArrayBox const* local_jy_ptr = local_jy[thread_num].get(); FArrayBox* global_jy_ptr = &cjyfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tby, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tby, thread_bx, { global_jy_ptr->atomicAdd(*local_jy_ptr, thread_bx, thread_bx, 0, 0, 1); }); FArrayBox const* local_jz_ptr = local_jz[thread_num].get(); FArrayBox* global_jz_ptr = &cjzfab; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbz, thread_bx, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(tbz, thread_bx, { global_jz_ptr->atomicAdd(*local_jz_ptr, thread_bx, thread_bx, 0, 0, 1); }); diff --git a/Source/WarpXEvolve.cpp b/Source/WarpXEvolve.cpp index ce6ed14d7..6983e1aba 100644 --- a/Source/WarpXEvolve.cpp +++ b/Source/WarpXEvolve.cpp @@ -468,7 +468,7 @@ WarpX::EvolveB (int lev, PatchType patch_type, amrex::Real dt) // Loop through the grids, and over the tiles within each grid #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*Bx, TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -506,7 +506,7 @@ WarpX::EvolveB (int lev, PatchType patch_type, amrex::Real dt) const auto& pml_E = (patch_type == PatchType::fine) ? pml[lev]->GetE_fp() : pml[lev]->GetE_cp(); #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*pml_B[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -593,7 +593,7 @@ WarpX::EvolveE (int lev, PatchType patch_type, amrex::Real dt) // Loop through the grids, and over the tiles within each grid #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*Ex, TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -650,7 +650,7 @@ WarpX::EvolveE (int lev, PatchType patch_type, amrex::Real dt) const auto& pml_E = (patch_type == PatchType::fine) ? pml[lev]->GetE_fp() : pml[lev]->GetE_cp(); const auto& pml_F = (patch_type == PatchType::fine) ? pml[lev]->GetF_fp() : pml[lev]->GetF_cp(); #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*pml_E[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -751,7 +751,7 @@ WarpX::EvolveF (int lev, PatchType patch_type, Real dt, DtType dt_type) const auto& pml_E = (patch_type == PatchType::fine) ? pml[lev]->GetE_fp() : pml[lev]->GetE_cp(); #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*pml_F, TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -797,7 +797,7 @@ WarpX::DampPML (int lev, PatchType patch_type) : pml[lev]->GetMultiSigmaBox_cp(); #ifdef _OPENMP -#pragma omp parallel if (!Cuda::inLaunchRegion()) +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*pml_E[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { diff --git a/Source/WarpXMove.cpp b/Source/WarpXMove.cpp index 2082130ef..c4eb2300d 100644 --- a/Source/WarpXMove.cpp +++ b/Source/WarpXMove.cpp @@ -224,7 +224,7 @@ WarpX::shiftMF (MultiFab& mf, const Geometry& geom, int num_shift, int dir) FArrayBox* srcfab = &(tmpmf[mfi]); Box outbox = mfi.fabbox(); outbox &= adjBox; - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(outbox, toutbox, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(outbox, toutbox, { srcfab->setVal(0.0, toutbox, 0, nc); }); @@ -234,7 +234,7 @@ WarpX::shiftMF (MultiFab& mf, const Geometry& geom, int num_shift, int dir) } else { dstBox.growLo(dir, num_shift); } - AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(dstBox, tdstBox, + AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA(dstBox, tdstBox, { dstfab->setVal(0.0, tdstBox, 0, nc); dstfab->copy(*srcfab, amrex::shift(tdstBox,dir,num_shift), 0, tdstBox, 0, nc); |