HDINLINE static DataSpace<DIM3> getBlockIndex(const Base &base,
const DataSpace<DIM3>& _blockIdx, uint32_t exchangeType)
{
DataSpace<DIM3> result(_blockIdx);
DataSpace<DIM3> directions = Mask::getRelativeDirections<DIM3 > (exchangeType);
if (directions.x() == 0)
result.x() += base.getGuardingSuperCells();
else
if (directions.x() == 1)
result.x() += base.getGridSuperCells().x() - base.getGuardingSuperCells();
if (directions.y() == 0)
result.y() += base.getGuardingSuperCells();
else
if (directions.y() == 1)
result.y() += base.getGridSuperCells().y() - base.getGuardingSuperCells();
if (directions.z() == 0)
result.z() += base.getGuardingSuperCells();
else
if (directions.z() == 1)
result.z() += base.getGridSuperCells().z() - base.getGuardingSuperCells();
return result;
}
void notify(uint32_t currentStep)
{
typedef typename MappingDesc::SuperCellSize SuperCellSize;
DataConnector& dc = Environment<>::get().DataConnector();
fieldE = &(dc.getData<FieldE > (FieldE::getName(), true));
fieldB = &(dc.getData<FieldB > (FieldB::getName(), true));
const int rank = Environment<simDim>::get().GridController().getGlobalRank();
getLineSliceFields < CORE + BORDER > ();
const SubGrid<simDim>& subGrid = Environment<simDim>::get().SubGrid();
// number of cells on the current CPU for each direction
const DataSpace<simDim> nrOfGpuCells = cellDescription->getGridLayout().getDataSpaceWithoutGuarding();
// global cell id offset (without guardings!)
// returns the global id offset of the "first" border cell on a GPU
const DataSpace<simDim> globalCellIdOffset(subGrid.getLocalDomain().offset);
// global number of cells for whole simulation: local cells on GPU * GPUs
// (assumed same size on each gpu :-/ -> todo: provide interface!)
//! \todo create a function for: global number of cells for whole simulation
//!
const DataSpace<simDim> globalNrOfCells = subGrid.getGlobalDomain().size;
/*FORMAT OUTPUT*/
/** \todo add float3_X with position of the cell to output*/
// check if the current GPU contains the "middle slice" along
// X_global / 2; Y_global / 2 over Z
if (globalCellIdOffset.x() <= globalNrOfCells.x() / 2 &&
globalCellIdOffset.x() + nrOfGpuCells.x() > globalNrOfCells.x() / 2)
#if(SIMDIM==DIM3)
if( globalCellIdOffset.z() <= globalNrOfCells.z() / 2 &&
globalCellIdOffset.z() + nrOfGpuCells.z() > globalNrOfCells.z() / 2)
#endif
for (int i = 0; i < nrOfGpuCells.y(); ++i)
{
const double xPos = double( i + globalCellIdOffset.y()) * SI::CELL_HEIGHT_SI;
outfile << currentStep << " " << rank << " ";
outfile << xPos << " "
/*<< sliceDataField->getHostBuffer().getDataBox()[i] */
<< double(sliceDataField->getHostBuffer().getDataBox()[i].x()) * UNIT_EFIELD << " "
<< double(sliceDataField->getHostBuffer().getDataBox()[i].y()) * UNIT_EFIELD << " "
<< double(sliceDataField->getHostBuffer().getDataBox()[i].z()) * UNIT_EFIELD << " "
<< "\n";
}
/* outfile << "[ANALYSIS] [" << rank << "] [COUNTER] [LineSliceFields] [" << currentStep << "] " <<
sliceDataField << "\n"; */
// free line to separate timesteps in gnuplot via the "index" option
outfile << std::endl;
}
static std::string dspToStr(DataSpace<DIM3>& dsp)
{
std::stringstream stream;
stream << "(" << dsp.x() << ", " << dsp.y() << ", " << dsp.z() << ")";
return stream.str();
}
HDINLINE static DataSpace<DIM3> getBlockIndex(const Base &base,
const DataSpace<DIM3>& _blockIdx, uint32_t exchangeType)
{
DataSpace<DIM3> result(_blockIdx);
DataSpace<DIM3> directions = Mask::getRelativeDirections<DIM3 > (exchangeType);
size_t guardingBlocks = base.getGuardingSuperCells();
size_t borderBlocks = base.getBorderSuperCells();
switch (directions.x())
{
case 0:
result.x() += guardingBlocks + borderBlocks;
break;
case -1:
result.x() += guardingBlocks;
break;
case 1:
result.x() += base.getGridSuperCells().x() - guardingBlocks -
borderBlocks;
break;
}
switch (directions.y())
{
case 0:
result.y() += guardingBlocks + borderBlocks;
break;
case -1:
result.y() += guardingBlocks;
break;
case 1:
result.y() += base.getGridSuperCells().y() - guardingBlocks -
borderBlocks;
break;
}
switch (directions.z())
{
case 0:
result.z() += guardingBlocks + borderBlocks;
break;
case -1:
result.z() += guardingBlocks;
break;
case 1:
result.z() += base.getGridSuperCells().z() - guardingBlocks -
borderBlocks;
break;
}
return result;
}
HINLINE static DataSpace<DIM3> getGridDim(const Base &base, const DataSpace<DIM3> &gBlocks)
{
const int x = gBlocks.x();
const int x_ = gBlocks.x() - 2 * base.getGuardingSuperCells();
const int y = gBlocks.y();
const int z_ = gBlocks.z() - 2 * base.getGuardingSuperCells();
return DataSpace<DIM3 > (x * y + z_ * y + x_*z_,
2 * base.getGuardingSuperCells(),
1);
}
HDINLINE static DataSpace<DIM3> getBlockIndex(const Base &base,
const DataSpace<DIM3> &gBlocks,
const DataSpace<DIM3>& _blockIdx)
{
const int g = base.getGuardingSuperCells();
const int b = base.getBorderSuperCells();
const int x = gBlocks.x() - 2 * g;
const int x_ = gBlocks.x() - 2 * g - 2 * b;
const int y = gBlocks.y() - 2 * g;
const int z = gBlocks.z() - 2 * g;
const int z_ = gBlocks.z() - 2 * g - 2 * b;
if (_blockIdx.x() < (x * y))
{
/* area is x*y */
const int tmp_x = _blockIdx.x();
return DataSpace<DIM3 > (tmp_x % x + g,
tmp_x / x + g,
g + _blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(z - b));
}
if ((_blockIdx.x() >= (x * y)) && _blockIdx.x() < (x * y + z_ * y))
{
/* area is z_*y */
const int tmp_x = _blockIdx.x() - (x * y);
return DataSpace<DIM3 > (g + _blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(x - b),
tmp_x / z_ + g,
(tmp_x % z_) + g + b);
}
/* area is x_*z_ */
const int tmp_x = _blockIdx.x() - (x * y) - (z_ * y);
return DataSpace<DIM3 > ((tmp_x % x_) + g + b,
g + _blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(y - b),
(tmp_x / x_) + g + b);
}
__global__ void kernelLineSliceFields(E_DataBox fieldE, B_DataBox fieldB,
float3_X* sliceDataField,
DataSpace<simDim> globalCellIdOffset,
DataSpace<simDim> globalNrOfCells,
Mapping mapper)
{
typedef typename Mapping::SuperCellSize SuperCellSize;
const DataSpace<simDim > threadIndex(threadIdx);
const int linearThreadIdx = DataSpaceOperations<simDim>::template map<SuperCellSize > (threadIndex);
const DataSpace<simDim> superCellIdx(mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx)));
__syncthreads();
// GPU-local cell id with lower GPU-local guarding
const DataSpace<simDim> localCell(superCellIdx * SuperCellSize::toRT() + threadIndex);
const float3_X b = fieldB(localCell);
const float3_X e = fieldE(localCell);
// GPU-local cell id without lower GPU-local guarding
const DataSpace<simDim> localCellWG(
localCell
- SuperCellSize::toRT() * mapper.getGuardingSuperCells());
// global cell id
const DataSpace<simDim> globalCell = localCellWG + globalCellIdOffset;
// slice out one cell along an axis
if ((globalCell.x() == globalNrOfCells.x() / 2))
#if(SIMDIM==DIM3)
if(globalCell.z() == globalNrOfCells.z() / 2)
#endif
sliceDataField[localCellWG.y()] = e;
__syncthreads();
}
/*! 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();
}
HDINLINE static DataSpace<DIM3> getBlockIndex(const Base &base,
const DataSpace<DIM3> &gBlocks,
const DataSpace<DIM3>& _blockIdx)
{
const int x = gBlocks.x();
const int x_ = gBlocks.x() - 2 * base.getGuardingSuperCells();
const int y = gBlocks.y();
const int z = gBlocks.z();
const int z_ = gBlocks.z() - 2 * base.getGuardingSuperCells();
if (_blockIdx.x() < (x * y))
{
/* area is x*y */
const int tmp_x = _blockIdx.x();
return DataSpace<DIM3 > (tmp_x % x,
tmp_x / x,
_blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(z - base.getGuardingSuperCells()));
}
if ((_blockIdx.x() >= (x * y)) && _blockIdx.x() < (x * y + z_ * y))
{
/* area is z_*y */
const int tmp_x = _blockIdx.x() - (x * y);
return DataSpace<DIM3 > (_blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(x - base.getGuardingSuperCells()),
tmp_x / z_,
(tmp_x % z_) + base.getGuardingSuperCells());
}
/* area is x_*z_ */
const int tmp_x = _blockIdx.x() - (x * y) - (z_ * y);
return DataSpace<DIM3 > ((tmp_x % x_) + base.getGuardingSuperCells(),
_blockIdx.y() / 2 +
(_blockIdx.y() & 1u) *
(y - base.getGuardingSuperCells()),
(tmp_x / x_) + base.getGuardingSuperCells());
}
HINLINE static DataSpace<DIM3> getGridDim(const Base &base, uint32_t exchangeType)
{
DataSpace<DIM3> result(base.getGridSuperCells() - 2 * base.getGuardingSuperCells());
DataSpace<DIM3> directions = Mask::getRelativeDirections<DIM3 > (exchangeType);
if (directions.x() != 0)
result.x() = base.getGuardingSuperCells();
if (directions.y() != 0)
result.y() = base.getGuardingSuperCells();
if (directions.z() != 0)
result.z() = base.getGuardingSuperCells();
return result;
}
HDINLINE RefValueType operator()(const DataSpace<DIM3> &idx = DataSpace<DIM3>())
{
return (Base::operator[](idx.z()))[idx.y()][idx.x()];
}
DINLINE void operator()(DataBoxJ dataBoxJ,
const PosType pos,
const VelType velocity,
const ChargeType charge,
const float_X deltaTime)
{
this->charge = charge;
const float3_X deltaPos = float3_X(velocity.x() * deltaTime / cellSize.x(),
velocity.y() * deltaTime / cellSize.y(),
velocity.z() * deltaTime / cellSize.z());
const PosType oldPos = pos - deltaPos;
Line<float3_X> line(oldPos, pos);
DataSpace<DIM3> gridShift;
/* Define in which direction the particle leaves the cell.
* It is not relevant whether the particle leaves the cell via
* the positive or negative cell border.
*
* 0 == stay in cell
* 1 == leave cell
*/
DataSpace<simDim> leaveCell;
/* calculate the offset for the virtual coordinate system */
for(int d=0; d<simDim; ++d)
{
int iStart;
int iEnd;
constexpr bool isSupportEven = ( supp % 2 == 0 );
RelayPoint< isSupportEven >()(
iStart,
iEnd,
line.m_pos0[d],
line.m_pos1[d]
);
gridShift[d] = iStart < iEnd ? iStart : iEnd; // integer min function
/* particle is leaving the cell */
leaveCell[d] = iStart != iEnd ? 1 : 0;
/* shift the particle position to the virtual coordinate system */
line.m_pos0[d] -= gridShift[d];
line.m_pos1[d] -= gridShift[d];
}
/* shift current field to the virtual coordinate system */
auto cursorJ = dataBoxJ.shift(gridShift).toCursor();
/**
* \brief the following three calls separate the 3D current deposition
* into three independent 1D calls, each for one direction and current component.
* Therefore the coordinate system has to be rotated so that the z-direction
* is always specific.
*/
using namespace cursor::tools;
cptCurrent1D(
DataSpace<simDim>(leaveCell.y(),leaveCell.z(),leaveCell.x()),
twistVectorFieldAxes<PMacc::math::CT::Int < 1, 2, 0 > >(cursorJ),
rotateOrigin < 1, 2, 0 > (line),
cellSize.x()
);
cptCurrent1D(
DataSpace<simDim>(leaveCell.z(),leaveCell.x(),leaveCell.y()),
twistVectorFieldAxes<PMacc::math::CT::Int < 2, 0, 1 > >(cursorJ),
rotateOrigin < 2, 0, 1 > (line),
cellSize.y()
);
cptCurrent1D(
leaveCell,
cursorJ,
line,
cellSize.z()
);
}
float3_64 getGlobalCell() const
{
return float3_64( typeCast<float_64>(globalCellOffset.x()) + typeCast<float_64>(position.x()),
typeCast<float_64>(globalCellOffset.y()) + typeCast<float_64>(position.y()),
typeCast<float_64>(globalCellOffset.z()) + typeCast<float_64>(position.z()) );
}
HDINLINE DataSpace<DIM3> reduce(const DataSpace<DIM3> &value)
{
z = value.z();
return DataSpace<DIM3 > (value.x() * z, value.y(), 1);
}
