The field ionization module seems to be failing on GPU when managed memory is turned o

<a class="issue-link js-issue-link" data-error-text="Failed to load title" data-id="16

With <a class="issue-link js-issue-link" data-error-text="Failed to load title" data-i

Managed Memory Leftovers: MR Init, QED, Ionization about warpx HOT 4 CLOSED

RemiLehe commented on July 30, 2024 1

Managed Memory Leftovers: MR Init, QED, Ionization

from warpx.

Comments (4)

ax3l commented on July 30, 2024

After looking at the logic, I think this affects QED, too.
Tested with ../../Examples/Tests/qed/quantum_synchrotron/inputs_2d
CC @lucafedeli88

Both use the same logic for creating new particles cc @atmyers

WarpX/Source/Particles/MultiParticleContainer.cpp

Lines 889 to 924 in 03bfb02

 auto& pc_product = allcontainers[pc_source->ionization_product]; 

 SmartCopyFactory copy_factory(*pc_source, *pc_product); 

 auto phys_pc_ptr = static_cast<PhysicalParticleContainer*>(pc_source.get()); 

 auto Copy = copy_factory.getSmartCopy(); 

 auto Transform = IonizationTransformFunc(); 

 pc_source ->defineAllParticleTiles(); 

 pc_product->defineAllParticleTiles(); 

 auto info = getMFItInfo(*pc_source, *pc_product); 

 #ifdef AMREX_USE_OMP 

 #pragma omp parallel if (Gpu::notInLaunchRegion()) 

 #endif 

 for (WarpXParIter pti(*pc_source, lev, info); pti.isValid(); ++pti) 

 { 

 if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) 

 { 

 amrex::Gpu::synchronize(); 

 } 

 Real wt = amrex::second(); 

 auto& src_tile = pc_source ->ParticlesAt(lev, pti); 

 auto& dst_tile = pc_product->ParticlesAt(lev, pti); 

 auto Filter = phys_pc_ptr->getIonizationFunc(pti, lev, Ex.nGrowVect(), 

 Ex[pti], Ey[pti], Ez[pti], 

 Bx[pti], By[pti], Bz[pti]); 

 const auto np_dst = dst_tile.numParticles(); 

 const auto num_added = filterCopyTransformParticles<1>(dst_tile, src_tile, np_dst, 

 Filter, Copy, Transform); 

 setNewParticleIDs(dst_tile, np_dst, num_added);

I see some page faults that are mitigated by managed memory with nsys, first DtoH and then also HtoD. Possible that the particle tiles of the source are device tiles, accessed without a copy from host, and that the dst tile is not actively copied back to device.

from warpx.

ax3l commented on July 30, 2024

#3825 is not yet released, removing respective tag.

Ionization fixed #3874
QED fix incoming via #3873

from warpx.

ax3l commented on July 30, 2024

Examples/Physics_applications/laser_acceleration/inputs_2d fails due to managed memory usage in WarpX::BuildBufferMasksInBox for all MR setups.

WarpX/Source/WarpX.cpp

Lines 2863 to 2918 in 03bfb02

 void 

 WarpX::BuildBufferMasksInBox ( const amrex::Box tbx, amrex::IArrayBox &buffer_mask, 

 const amrex::IArrayBox &guard_mask, const int ng ) 

 { 

 bool setnull; 

 const amrex::Dim3 lo = amrex::lbound( tbx ); 

 const amrex::Dim3 hi = amrex::ubound( tbx ); 

 Array4<int> msk = buffer_mask.array(); 

 Array4<int const> gmsk = guard_mask.array(); 

 #if defined(WARPX_DIM_1D_Z) 

 int k = lo.z; 

 int j = lo.y; 

 for (int i = lo.x; i <= hi.x; ++i) { 

 setnull = false; 

 // If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region 

 for (int ii = i-ng; ii <= i+ng; ++ii) { 

 if ( gmsk(ii,j,k) == 0 ) setnull = true; 

 } 

 if ( setnull ) msk(i,j,k) = 0; 

 else msk(i,j,k) = 1; 

 } 

 #elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) 

 int k = lo.z; 

 for (int j = lo.y; j <= hi.y; ++j) { 

 for (int i = lo.x; i <= hi.x; ++i) { 

 setnull = false; 

 // If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region 

 for (int jj = j-ng; jj <= j+ng; ++jj) { 

 for (int ii = i-ng; ii <= i+ng; ++ii) { 

 if ( gmsk(ii,jj,k) == 0 ) setnull = true; 

 } 

 } 

 if ( setnull ) msk(i,j,k) = 0; 

 else msk(i,j,k) = 1; 

 } 

 } 

 #elif defined(WARPX_DIM_3D) 

 for (int k = lo.z; k <= hi.z; ++k) { 

 for (int j = lo.y; j <= hi.y; ++j) { 

 for (int i = lo.x; i <= hi.x; ++i) { 

 setnull = false; 

 // If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region 

 for (int kk = k-ng; kk <= k+ng; ++kk) { 

 for (int jj = j-ng; jj <= j+ng; ++jj) { 

 for (int ii = i-ng; ii <= i+ng; ++ii) { 

 if ( gmsk(ii,jj,kk) == 0 ) setnull = true; 

 } 

 } 

 } 

 if ( setnull ) msk(i,j,k) = 0; 

 else msk(i,j,k) = 1; 

 } 

 } 

 } 

 #endif 

 }

It's consequently also very slow. To fix, we need an update using ParallelFor on the right box extents.

cc @WeiqunZhang @EZoni

from warpx.

ax3l commented on July 30, 2024

With #3880 merged, we have all known issues fixed again :)

Thanks everyone that reported, tested & helped! 🤝

from warpx.

Managed Memory Leftovers: MR Init, QED, Ionization about warpx HOT 4 CLOSED

Comments (4)

Related Issues (20)

Recommend Projects

React

Vue.js

Typescript

TensorFlow

Django

Laravel

D3

Recommend Topics

javascript

web

server

Machine learning

Visualization

Game

Recommend Org

Facebook

Microsoft

Google

Alibaba

D3

Tencent

	auto& pc_product = allcontainers[pc_source->ionization_product];

	SmartCopyFactory copy_factory(pc_source, pc_product);
	auto phys_pc_ptr = static_cast<PhysicalParticleContainer*>(pc_source.get());

	auto Copy = copy_factory.getSmartCopy();
	auto Transform = IonizationTransformFunc();

	pc_source ->defineAllParticleTiles();
	pc_product->defineAllParticleTiles();

	auto info = getMFItInfo(pc_source, pc_product);

	#ifdef AMREX_USE_OMP
	#pragma omp parallel if (Gpu::notInLaunchRegion())
	#endif
	for (WarpXParIter pti(*pc_source, lev, info); pti.isValid(); ++pti)
	{
	if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
	{
	amrex::Gpu::synchronize();
	}
	Real wt = amrex::second();

	auto& src_tile = pc_source ->ParticlesAt(lev, pti);
	auto& dst_tile = pc_product->ParticlesAt(lev, pti);

	auto Filter = phys_pc_ptr->getIonizationFunc(pti, lev, Ex.nGrowVect(),
	Ex[pti], Ey[pti], Ez[pti],
	Bx[pti], By[pti], Bz[pti]);

	const auto np_dst = dst_tile.numParticles();
	const auto num_added = filterCopyTransformParticles<1>(dst_tile, src_tile, np_dst,
	Filter, Copy, Transform);

	setNewParticleIDs(dst_tile, np_dst, num_added);

	void
	WarpX::BuildBufferMasksInBox ( const amrex::Box tbx, amrex::IArrayBox &buffer_mask,
	const amrex::IArrayBox &guard_mask, const int ng )
	{
	bool setnull;
	const amrex::Dim3 lo = amrex::lbound( tbx );
	const amrex::Dim3 hi = amrex::ubound( tbx );
	Array4<int> msk = buffer_mask.array();
	Array4<int const> gmsk = guard_mask.array();
	#if defined(WARPX_DIM_1D_Z)
	int k = lo.z;
	int j = lo.y;
	for (int i = lo.x; i <= hi.x; ++i) {
	setnull = false;
	// If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region
	for (int ii = i-ng; ii <= i+ng; ++ii) {
	if ( gmsk(ii,j,k) == 0 ) setnull = true;
	}
	if ( setnull ) msk(i,j,k) = 0;
	else msk(i,j,k) = 1;
	}
	#elif defined(WARPX_DIM_XZ) \|\| defined(WARPX_DIM_RZ)
	int k = lo.z;
	for (int j = lo.y; j <= hi.y; ++j) {
	for (int i = lo.x; i <= hi.x; ++i) {
	setnull = false;
	// If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region
	for (int jj = j-ng; jj <= j+ng; ++jj) {
	for (int ii = i-ng; ii <= i+ng; ++ii) {
	if ( gmsk(ii,jj,k) == 0 ) setnull = true;
	}
	}
	if ( setnull ) msk(i,j,k) = 0;
	else msk(i,j,k) = 1;
	}
	}
	#elif defined(WARPX_DIM_3D)
	for (int k = lo.z; k <= hi.z; ++k) {
	for (int j = lo.y; j <= hi.y; ++j) {
	for (int i = lo.x; i <= hi.x; ++i) {
	setnull = false;
	// If gmsk=0 for any neighbor within ng cells, current cell is in the buffer region
	for (int kk = k-ng; kk <= k+ng; ++kk) {
	for (int jj = j-ng; jj <= j+ng; ++jj) {
	for (int ii = i-ng; ii <= i+ng; ++ii) {
	if ( gmsk(ii,jj,kk) == 0 ) setnull = true;
	}
	}
	}
	if ( setnull ) msk(i,j,k) = 0;
	else msk(i,j,k) = 1;
	}
	}
	}
	#endif
	}