From c0648b9d015f17cf8c9748480b5a1d0c1943a41c Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 23 Oct 2018 12:54:03 -0400 Subject: Revert "Launch all of these in one kernel (from Kevin Gott)" This reverts commit 348cb47479a307cdae2860ba993d22ba172fb823. --- Source/PhysicalParticleContainer.cpp | 57 +++++------------------------------- 1 file changed, 7 insertions(+), 50 deletions(-) (limited to 'Source/PhysicalParticleContainer.cpp') diff --git a/Source/PhysicalParticleContainer.cpp b/Source/PhysicalParticleContainer.cpp index 1fe9cff0e..b9fb45f23 100644 --- a/Source/PhysicalParticleContainer.cpp +++ b/Source/PhysicalParticleContainer.cpp @@ -725,12 +725,10 @@ PhysicalParticleContainer::Evolve (int lev, #pragma omp parallel #endif { - Cuda::DeviceVector xp, yp, zp, giv; - - std::unique_ptr local_rho(new FArrayBox()); - std::unique_ptr local_jx(new FArrayBox()); - std::unique_ptr local_jy(new FArrayBox()); - std::unique_ptr local_jz(new FArrayBox()); + if (local_rho == nullptr) local_rho.reset(new amrex::FArrayBox()); + if (local_jx == nullptr) local_jx.reset( new amrex::FArrayBox()); + if (local_jy == nullptr) local_jy.reset( new amrex::FArrayBox()); + if (local_jz == nullptr) local_jz.reset( new amrex::FArrayBox()); FArrayBox filtered_Ex, filtered_Ey, filtered_Ez; FArrayBox filtered_Bx, filtered_By, filtered_Bz; @@ -1218,16 +1216,6 @@ PhysicalParticleContainer::Evolve (int lev, jz_ptr = local_jz->dataPtr(); FArrayBox* local_jx_ptr = local_jx.get(); - FArrayBox* local_jy_ptr = local_jy.get(); - FArrayBox* local_jz_ptr = local_jz.get(); - - AMREX_CUDA_LAUNCH_HOST_DEVICE_XYZ( tbx, tby, tbz, bx, by, bz, - { local_jx_ptr->setVal(0.0, bx, 0, 1); }, - { local_jy_ptr->setVal(0.0, by, 0, 1); }, - { local_jz_ptr->setVal(0.0, bz, 0, 1); } - ); - -/* AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, b, { local_jx_ptr->setVal(0.0, b, 0, 1); @@ -1244,7 +1232,7 @@ PhysicalParticleContainer::Evolve (int lev, { local_jz_ptr->setVal(0.0, b, 0, 1); }); -*/ + jxntot = local_jx->length(); jyntot = local_jy->length(); jzntot = local_jz->length(); @@ -1269,22 +1257,6 @@ PhysicalParticleContainer::Evolve (int lev, BL_PROFILE_VAR_START(blp_accumulate); - FArrayBox const* local_jx_const_ptr = local_jx.get(); - FArrayBox const* local_jy_const_ptr = local_jy.get(); - FArrayBox const* local_jz_const_ptr = local_jz.get(); - - FArrayBox* global_jx_ptr = &jxfab; - FArrayBox* global_jy_ptr = &jyfab; - FArrayBox* global_jz_ptr = &jzfab; - - AMREX_CUDA_LAUNCH_HOST_DEVICE_XYZ(tbx,tby,tbz,thread_bx,thread_by,thread_bz, - { global_jx_ptr->atomicAdd(*local_jx_const_ptr, thread_bx, thread_bx, 0, 0, 1); }, - { global_jy_ptr->atomicAdd(*local_jy_const_ptr, thread_by, thread_by, 0, 0, 1); }, - { global_jz_ptr->atomicAdd(*local_jz_const_ptr, thread_bz, thread_bz, 0, 0, 1); } - ); - - -/* FArrayBox const* local_jx_const_ptr = local_jx.get(); FArrayBox* global_jx_ptr = &jxfab; AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, @@ -1305,7 +1277,7 @@ PhysicalParticleContainer::Evolve (int lev, { global_jz_ptr->atomicAdd(*local_jz_const_ptr, thread_bx, thread_bx, 0, 0, 1); }); -*/ + BL_PROFILE_VAR_STOP(blp_accumulate); } @@ -1363,21 +1335,6 @@ PhysicalParticleContainer::Evolve (int lev, BL_PROFILE_VAR_START(blp_accumulate); - FArrayBox const* local_jx_ptr = local_jx.get(); - FArrayBox const* local_jy_ptr = local_jy.get(); - FArrayBox const* local_jz_ptr = local_jz.get(); - - FArrayBox* global_jx_ptr = &cjxfab; - FArrayBox* global_jy_ptr = &cjyfab; - FArrayBox* global_jz_ptr = &cjzfab; - - AMREX_CUDA_LAUNCH_HOST_DEVICE_XYZ(tbx, tby, tbz, thread_bx, thread_by, thread_bz, - { global_jx_ptr->atomicAdd(*local_jx_ptr, thread_bx, thread_bx, 0, 0, 1); }, - { global_jy_ptr->atomicAdd(*local_jy_ptr, thread_by, thread_by, 0, 0, 1); }, - { global_jz_ptr->atomicAdd(*local_jz_ptr, thread_bz, thread_bz, 0, 0, 1); } - ); - -/* FArrayBox const* local_jx_ptr = local_jx.get(); FArrayBox* global_jx_ptr = &cjxfab; AMREX_CUDA_LAUNCH_HOST_DEVICE_LAMBDA(tbx, thread_bx, @@ -1398,7 +1355,7 @@ PhysicalParticleContainer::Evolve (int lev, { global_jz_ptr->atomicAdd(*local_jz_ptr, thread_bx, thread_bx, 0, 0, 1); }); -*/ + BL_PROFILE_VAR_STOP(blp_accumulate); } -- cgit v1.2.3