aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGravatar Weiqun Zhang <weiqunzhang@lbl.gov> 2018-10-26 10:48:23 -0700
committerGravatar Weiqun Zhang <weiqunzhang@lbl.gov> 2018-10-26 10:48:23 -0700
commitf14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c (patch)
tree78f1cc8ae3ed551d4b3dd76b72e589d42561eb46
parent8642045c0f27a087430a0a1083db7fb678e8fe2a (diff)
downloadWarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.tar.gz
WarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.tar.zst
WarpX-f14a9bb95b2b9f7bb71b06daf7ce7e4ad1ff0c7c.zip
update according to amrex changes
-rw-r--r--Source/PhysicalParticleContainer.H2
-rw-r--r--Source/PhysicalParticleContainer.cpp26
-rw-r--r--Source/WarpXEvolve.cpp12
-rw-r--r--Source/WarpXMove.cpp4
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);