/** constructor
*
* @param size extent for each dimension (in elements)
*/
MappedBufferIntern(DataSpace<DIM> size):
DeviceBuffer<TYPE, DIM>(size, size),
pointer(nullptr), ownPointer(true)
{
#if( PMACC_CUDA_ENABLED == 1 )
CUDA_CHECK((cuplaError_t)cudaHostAlloc(&pointer, size.productOfComponents() * sizeof (TYPE), cudaHostAllocMapped));
#else
pointer = new TYPE[size.productOfComponents()];
#endif
reset(false);
}
virtual void init()
{
// __startAtomicTransaction(__getTransactionEvent());
size_t current_size = host->getCurrentSize();
DataSpace<DIM> hostCurrentSize = host->getCurrentDataSpace(current_size);
if (host->is1D() && device->is1D())
fastCopy(host->getPointer(), device->getPointer(), hostCurrentSize.productOfComponents());
else
copy(hostCurrentSize);
device->setCurrentSize(current_size);
this->activate();
// __setTransactionEvent(__endTransaction());
}
/**
* Add Exchange in dedicated memory space.
*
* An Exchange is added to this GridBuffer. The exchange buffers use
* the their own memory instead of using the GridBuffer's memory space.
*
* @param receive a Mask which describes the directions for the exchange
* @param dataSpace size of the newly created exchange buffer in each dimension
* @param sizeOnDevice if true, internal buffers have their size information on the device, too
*/
void addExchangeBuffer(const Mask &receive, const DataSpace<DIM> &dataSpace, uint32_t communicationTag, bool sizeOnDevice = false)
{
if (hasOneExchange && (communicationTag != lastUsedCommunicationTag))
throw std::runtime_error("It is not allowed to give the same GridBuffer different communicationTags");
lastUsedCommunicationTag = communicationTag;
/*don't create buffer with 0 (zero) elements*/
if (dataSpace.productOfComponents() != 0)
{
receiveMask = receiveMask + receive;
sendMask = this->receiveMask.getMirroredMask();
Mask send = receive.getMirroredMask();
for (uint32_t ex = 1; ex < 27; ++ex)
{
if (send.isSet(ex))
{
uint32_t uniqCommunicationTag = (communicationTag << 5) | ex;
if (!hasOneExchange && !privateGridBuffer::UniquTag::getInstance().isTagUniqu(uniqCommunicationTag))
{
std::stringstream message;
message << "unique exchange communication tag ("
<< uniqCommunicationTag << ") witch is created from communicationTag ("
<< communicationTag << ") allready used for other gridbuffer exchange";
throw std::runtime_error(message.str());
}
hasOneExchange = true;
if (sendExchanges[ex] != NULL)
{
throw std::runtime_error("Exchange already added!");
}
//GridLayout<DIM> memoryLayout(size);
maxExchange = std::max(maxExchange, ex + 1u);
sendExchanges[ex] = new ExchangeIntern<BORDERTYPE, DIM > (/*memoryLayout*/ dataSpace,
ex, uniqCommunicationTag, sizeOnDevice);
ExchangeType recvex = Mask::getMirroredExchangeType(ex);
maxExchange = std::max(maxExchange, recvex + 1u);
receiveExchanges[recvex] = new ExchangeIntern<BORDERTYPE, DIM > (/*memoryLayout*/ dataSpace,
recvex, uniqCommunicationTag, sizeOnDevice);
}
}
}
}
/*! initializes all processes to build a 3D-grid
*
* @param nodes number of GPU nodes in each dimension
* @param periodic specifying whether the grid is periodic (1) or not (0) in each dimension
*
* \warning throws invalid argument if cx*cy*cz != totalnodes
*/
void init(DataSpace<DIM3> numberProcesses, DataSpace<DIM3> periodic) throw (std::invalid_argument)
{
this->periodic = periodic;
//check if parameters are correct
MPI_CHECK(MPI_Comm_size(MPI_COMM_WORLD, &mpiSize));
if (numberProcesses.productOfComponents() != mpiSize)
{
throw std::invalid_argument("wrong parameters or wrong mpirun-call!");
}
//1. create Communicator (computing_comm) of computing nodes (ranks 0...n)
MPI_Comm computing_comm = MPI_COMM_WORLD;
yoffset = 0;
// 2. create topology
//int dims[3];
dims[0] = numberProcesses.x();
dims[1] = numberProcesses.y();
dims[2] = numberProcesses.z();
topology = MPI_COMM_NULL;
int periods[] = {periodic.x(), periodic.y(), periodic.z()};
/*create new communicator based on cartesian coordinates*/
MPI_CHECK(MPI_Cart_create(computing_comm, DIM, dims, periods, 0, &topology));
// 3. update Host rank
hostRank = UpdateHostRank();
//4. update Coordinates
updateCoordinates();
}
DINLINE void
ComputeGridValuePerFrame<T_ParticleShape, T_DerivedAttribute>::operator()
(
T_Acc const & acc,
FrameType& frame,
const int localIdx,
const TVecSuperCell superCell,
BoxTmp& tmpBox
)
{
/* \todo in the future and if useful, the functor can be a parameter */
T_DerivedAttribute particleAttribute;
auto particle = frame[localIdx];
/* particle attributes: in-cell position and generic, derived attribute */
const floatD_X pos = particle[position_];
const auto particleAttr = particleAttribute( particle );
/** Shift to the cell the particle belongs to
* range of particleCell: [DataSpace<simDim>::create(0), TVecSuperCell]
*/
const int particleCellIdx = particle[localCellIdx_];
const DataSpace<TVecSuperCell::dim> particleCell(
DataSpaceOperations<TVecSuperCell::dim>::map( superCell, particleCellIdx )
);
auto fieldTmpShiftToParticle = tmpBox.shift(particleCell);
/* loop around the particle's cell (according to shape) */
const DataSpace<simDim> lowMargin(LowerMargin().toRT());
const DataSpace<simDim> upMargin(UpperMargin().toRT());
const DataSpace<simDim> marginSpace(upMargin + lowMargin + 1);
const int numWriteCells = marginSpace.productOfComponents();
for (int i = 0; i < numWriteCells; ++i)
{
/** for the current cell i the multi dimensional index currentCell is only positive:
* allowed range = [DataSpace<simDim>::create(0), LowerMargin+UpperMargin]
*/
const DataSpace<simDim> currentCell = DataSpaceOperations<simDim>::map(marginSpace, i);
/** calculate the offset between the current cell i with simDim index currentCell
* and the cell of the particle (particleCell) in cells
*/
const DataSpace<simDim> offsetParticleCellToCurrentCell = currentCell - lowMargin;
/** assign particle contribution component-wise to the lower left corner of
* the cell i
* \todo take care of non-yee cells
*/
float_X assign( 1.0 );
for (uint32_t d = 0; d < simDim; ++d)
assign *= AssignmentFunction()(float_X(offsetParticleCellToCurrentCell[d]) - pos[d]);
/** add contribution of the particle times the generic attribute
* to cell i
* note: the .x() is used because FieldTmp is a scalar field with only
* one "x" component
*/
atomicAdd(
&(fieldTmpShiftToParticle(offsetParticleCellToCurrentCell).x()),
assign * particleAttr,
::alpaka::hierarchy::Threads{}
);
}
}
/**
* constructor
* @param dataSpace description of spread of any dimension
*/
Buffer(DataSpace<DIM> dataSpace) :
data_space(dataSpace),data1D(true)
{
CUDA_CHECK(cudaMallocHost(¤t_size, sizeof (size_t)));
*current_size = dataSpace.productOfComponents();
}
/** constructor
*
* @param size extent for each dimension (in elements)
* if the buffer is a view to an existing buffer the size
* can be less than `physicalMemorySize`
* @param physicalMemorySize size of the physical memory (in elements)
*/
Buffer(DataSpace<DIM> size, DataSpace<DIM> physicalMemorySize) :
data_space(size), data1D(true), current_size(NULL), m_physicalMemorySize(physicalMemorySize)
{
CUDA_CHECK(cudaMallocHost(¤t_size, sizeof (size_t)));
*current_size = size.productOfComponents();
}
/**
* Constructor
* \see GridBuffer
*/
HeapBuffer(DataSpace<DIM1> dataSpace) :
GridBuffer<VALUE, DIM1, BORDERVALUE>(dataSpace), ringBuffer(NULL)
{
ringBuffer = new RingBuffer<TYPE, TYPE > (dataSpace.productOfComponents());
}
