C++ Mapping::getGuardingSuperCells方法代码示例

__global__ void kernelPositionsParticles(ParticlesBox<FRAME, simDim> pb,
                                         SglParticle<FloatPos>* gParticle,
                                         Mapping mapper)
{

    __shared__ FRAME *frame;
    __shared__ bool isValid;
    __syncthreads(); /*wait that all shared memory is initialised*/

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

    if (linearThreadIdx == 0)
    {
        frame = &(pb.getLastFrame(superCellIdx, isValid));
    }

    __syncthreads();
    if (!isValid)
        return; //end kernel if we have no frames

    bool isParticle = (*frame)[linearThreadIdx][multiMask_];

    while (isValid)
    {
        if (isParticle)
        {
            PMACC_AUTO(particle,(*frame)[linearThreadIdx]);
            gParticle->position = particle[position_];
            gParticle->momentum = particle[momentum_];
            gParticle->weighting = particle[weighting_];
            gParticle->mass = frame->getMass(gParticle->weighting);
            gParticle->charge = frame->getCharge(gParticle->weighting);
            gParticle->gamma = Gamma<>()(gParticle->momentum, gParticle->mass);

            // storage number in the actual frame
            const lcellId_t frameCellNr = particle[localCellIdx_];

            // offset in the actual superCell = cell offset in the supercell
            const DataSpace<simDim> frameCellOffset(DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (frameCellNr));


            gParticle->globalCellOffset = (superCellIdx - mapper.getGuardingSuperCells())
                * MappingDesc::SuperCellSize::getDataSpace()
                + frameCellOffset;
        }
        __syncthreads();
        if (linearThreadIdx == 0)
        {
            frame = &(pb.getPreviousFrame(*frame, isValid));
        }
        isParticle = true;
        __syncthreads();
    }

}

__global__ void kernelCountParticles(PBox pb,
                                     uint64_cu* gCounter,
                                     Filter filter,
                                     Mapping mapper)
{

    typedef typename PBox::FrameType FRAME;
    const uint32_t Dim = Mapping::Dim;

    __shared__ FRAME *frame;
    __shared__ bool isValid;
    __shared__ int counter;
    __shared__ lcellId_t particlesInSuperCell;


    __syncthreads(); /*wait that all shared memory is initialised*/

    typedef typename Mapping::SuperCellSize SuperCellSize;

    const DataSpace<Dim > threadIndex(threadIdx);
    const int linearThreadIdx = DataSpaceOperations<Dim>::template map<SuperCellSize > (threadIndex);
    const DataSpace<Dim> superCellIdx(mapper.getSuperCellIndex(DataSpace<Dim > (blockIdx)));

    if (linearThreadIdx == 0)
    {
        frame = &(pb.getLastFrame(superCellIdx, isValid));
        particlesInSuperCell = pb.getSuperCell(superCellIdx).getSizeLastFrame();
        counter = 0;
    }
    __syncthreads();
    if (!isValid)
        return; //end kernel if we have no frames
    filter.setSuperCellPosition((superCellIdx - mapper.getGuardingSuperCells()) * mapper.getSuperCellSize());
    while (isValid)
    {
        if (linearThreadIdx < particlesInSuperCell)
        {
            if (filter(*frame, linearThreadIdx))
                atomicAdd(&counter, 1);
        }
        __syncthreads();
        if (linearThreadIdx == 0)
        {
            frame = &(pb.getPreviousFrame(*frame, isValid));
            particlesInSuperCell = math::CT::volume<SuperCellSize>::type::value;
        }
        __syncthreads();
    }

    __syncthreads();
    if (linearThreadIdx == 0)
    {
        atomicAdd(gCounter, (uint64_cu) counter);
    }
}

__global__ void CountMakroParticle(ParBox parBox, CounterBox counterBox, Mapping mapper)
{

    typedef MappingDesc::SuperCellSize SuperCellSize;
    typedef typename ParBox::FrameType FrameType;

    const DataSpace<simDim> block(mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx)));
    /* counterBox has no guarding supercells*/
    const DataSpace<simDim> counterCell = block - mapper.getGuardingSuperCells();

    const DataSpace<simDim > threadIndex(threadIdx);
    const int linearThreadIdx = DataSpaceOperations<simDim>::template map<SuperCellSize > (threadIndex);

    __shared__ uint64_cu counterValue;
    __shared__ FrameType *frame;
    __shared__ bool isValid;

    if (linearThreadIdx == 0)
    {
        counterValue = 0;
        frame = &(parBox.getLastFrame(block, isValid));
        if (!isValid)
        {
            counterBox(counterCell) = counterValue;
        }
    }
    __syncthreads();
    if (!isValid)
        return; //end kernel if we have no frames

    bool isParticle = (*frame)[linearThreadIdx][multiMask_];

    while (isValid)
    {
        if (isParticle)
        {
            atomicAdd(&counterValue, static_cast<uint64_cu> (1LU));
        }
        __syncthreads();
        if (linearThreadIdx == 0)
        {
            frame = &(parBox.getPreviousFrame(*frame, isValid));
        }
        isParticle = true;
        __syncthreads();
    }

    if (linearThreadIdx == 0)
        counterBox(counterCell) = counterValue;
}

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

__global__ void kernelPaintFields(
                                  EBox fieldE,
                                  BBox fieldB,
                                  JBox fieldJ,
                                  DataBox<PitchedBox<float3_X, DIM2> > image,
                                  DataSpace<DIM2> transpose,
                                  const int slice,
                                  const uint32_t globalOffset,
                                  const uint32_t sliceDim,
                                  Mapping mapper)
{
    typedef typename MappingDesc::SuperCellSize Block;
    const DataSpace<simDim> threadId(threadIdx);
    const DataSpace<simDim> block = mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx));
    const DataSpace<simDim> cell(block * Block::getDataSpace() + threadId);
    const DataSpace<simDim> blockOffset((block - mapper.getGuardingSuperCells()) * Block::getDataSpace());


    const DataSpace<simDim> realCell(cell - MappingDesc::SuperCellSize::getDataSpace() * mapper.getGuardingSuperCells()); //delete guard from cell idx
    const DataSpace<DIM2> imageCell(
                                    realCell[transpose.x()],
                                    realCell[transpose.y()]);
    const DataSpace<simDim> realCell2(blockOffset + threadId); //delete guard from cell idx

#if (SIMDIM==DIM3)
    uint32_t globalCell = realCell2[sliceDim] + globalOffset;

    if (globalCell != slice)
        return;
#endif
    // set fields of this cell to vars
    typename BBox::ValueType field_b = fieldB(cell);
    typename EBox::ValueType field_e = fieldE(cell);
    typename JBox::ValueType field_j = fieldJ(cell);

#if(SIMDIM==DIM3)
    field_j = float3_X(
                       field_j.x() * CELL_HEIGHT * CELL_DEPTH,
                       field_j.y() * CELL_WIDTH * CELL_DEPTH,
                       field_j.z() * CELL_WIDTH * CELL_HEIGHT
                       );
#elif (SIMDIM==DIM2)
    field_j = float3_X(
                       field_j.x() * CELL_HEIGHT,
                       field_j.y() * CELL_WIDTH,
                       field_j.z() * CELL_WIDTH * CELL_HEIGHT
                       );
#endif

    // reset picture to black
    //   color range for each RGB channel: [0.0, 1.0]
    float3_X pic = float3_X(0., 0., 0.);

    // typical values of the fields to normalize them to [0,1]
    //
    pic.x() = visPreview::preChannel1(field_b / typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().x(),
                                      field_e / typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().y(),
                                      field_j / typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().z());
    pic.y() = visPreview::preChannel2(field_b / typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().x(),
                                      field_e / typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().y(),
                                      field_j / typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().z());
    pic.z() = visPreview::preChannel3(field_b / typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().x(),
                                      field_e / typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().y(),
                                      field_j / typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().z());
    //visPreview::preChannel1Col::addRGB(pic,
    //                                   visPreview::preChannel1(field_b * typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().x(),
    //                                                           field_e * typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().y(),
    //                                                           field_j * typicalFields<EM_FIELD_SCALE_CHANNEL1>::get().z()),
    //                                   visPreview::preChannel1_opacity);
    //visPreview::preChannel2Col::addRGB(pic,
    //                                   visPreview::preChannel2(field_b * typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().x(),
    //                                                           field_e * typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().y(),
    //                                                           field_j * typicalFields<EM_FIELD_SCALE_CHANNEL2>::get().z()),
    //                                   visPreview::preChannel2_opacity);
    //visPreview::preChannel3Col::addRGB(pic,
    //                                   visPreview::preChannel3(field_b * typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().x(),
    //                                                           field_e * typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().y(),
    //                                                           field_j * typicalFields<EM_FIELD_SCALE_CHANNEL3>::get().z()),
    //                                   visPreview::preChannel3_opacity);


    // draw to (perhaps smaller) image cell
    image(imageCell) = pic;
}

    DINLINE void operator()(ParBoxIons ionBox,
                                          ParBoxElectrons electronBox,
                                          FrameIonizer frameIonizer,
                                          Mapping mapper) const
    {

        /* "particle box" : container/iterator where the particles live in
         * and where one can get the frame in a super cell from
         */
        typedef typename ParBoxElectrons::FrameType ELECTRONFRAME;
        typedef typename ParBoxIons::FrameType IONFRAME;
        typedef typename ParBoxIons::FramePtr IonFramePtr;
        typedef typename ParBoxElectrons::FramePtr ElectronFramePtr;

        /* specify field to particle interpolation scheme */
        typedef typename PMacc::traits::Resolve<
            typename GetFlagType<IONFRAME,interpolation<> >::type
        >::type InterpolationScheme;

        /* margins around the supercell for the interpolation of the field on the cells */
        typedef typename GetMargin<InterpolationScheme>::LowerMargin LowerMargin;
        typedef typename GetMargin<InterpolationScheme>::UpperMargin UpperMargin;

        /* relevant area of a block */
        typedef SuperCellDescription<
            typename MappingDesc::SuperCellSize,
            LowerMargin,
            UpperMargin
            > BlockDescription_;

        /* for not mixing operations::assign up with the nvidia functor assign */
        namespace partOp = PMacc::particles::operations;

        /* definitions for domain variables, like indices of blocks and threads */
        typedef typename BlockDescription_::SuperCellSize SuperCellSize;
        /* multi-dimensional offset vector from local domain origin on GPU in units of super cells */
        const DataSpace<simDim> block(mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx)));

        /* multi-dim vector from origin of the block to a cell in units of cells */
        const DataSpace<simDim > threadIndex(threadIdx);
        /* conversion from a multi-dim cell coordinate to a linear coordinate of the cell in its super cell */
        const int linearThreadIdx = DataSpaceOperations<simDim>::template map<SuperCellSize > (threadIndex);

        /* multi-dim offset from the origin of the local domain on GPU
         * to the origin of the block of the in unit of cells
         */
        const DataSpace<simDim> blockCell = block * SuperCellSize::toRT();

        /* subtract guarding cells to only have the simulation volume */
        const DataSpace<simDim> localCellIndex = (block * SuperCellSize::toRT() + threadIndex) - mapper.getGuardingSuperCells() * SuperCellSize::toRT();

        /* typedef for the functor that writes new macro electrons into electron frames during runtime */
        typedef typename particles::ionization::WriteElectronIntoFrame WriteElectronIntoFrame;


        PMACC_SMEM( ionFrame, IonFramePtr );
        PMACC_SMEM( electronFrame,ElectronFramePtr );
        PMACC_SMEM( maxParticlesInFrame, lcellId_t );


        /* find last frame in super cell
         * define maxParticlesInFrame as the maximum frame size
         */
        if (linearThreadIdx == 0)
        {
            ionFrame = ionBox.getLastFrame(block);
            maxParticlesInFrame = PMacc::math::CT::volume<SuperCellSize>::type::value;
        }

        __syncthreads();
        if (!ionFrame.isValid())
            return; //end kernel if we have no frames

        /* caching of E- and B- fields and initialization of random generator if needed */
        frameIonizer.init(blockCell, linearThreadIdx, localCellIndex);

        /* Declare counter in shared memory that will later tell the current fill level or
         * occupation of the newly created target electron frames.
         */
        PMACC_SMEM( newFrameFillLvl, int );

        /* Declare local variable oldFrameFillLvl for each thread */
        int oldFrameFillLvl;

        /* Initialize local (register) counter for each thread
         * - describes how many new macro electrons should be created
         */
        unsigned int newMacroElectrons = 0;

        /* Declare local electron ID
         * - describes at which position in the new frame the new electron is to be created
         */
        int electronId;

        /* Master initializes the frame fill level with 0 */
        if (linearThreadIdx == 0)
        {
            newFrameFillLvl = 0;
            electronFrame = nullptr;
        }
//.........这里部分代码省略.........