C++ DataSpace::z方法代码示例

        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();
        }

        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 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);
 }