hoomd/mpcd/SorterGPU.cc
/*!
 * The sorted order is applied by swapping out the alternate per-particle data
 * arrays. The communication flags are \b not sorted in MPI because by design,
 * the caller is responsible for clearing out any old flags before using them.
 */
The sorted order is applied by swapping out the alternate per-particle data
arrays. The communication flags are \b not sorted in MPI because by design,
the caller is responsible for clearing out any old flags before using them.
The sorted order is applied by swapping out the alternate per - particle data arrays . The communication flags are \ b not sorted in MPI because by design the caller is responsible for clearing out any old flags before using them .
void mpcd::SorterGPU::applyOrder() const
    {
        {
        ArrayHandle<unsigned int> d_order(m_order, access_location::device, access_mode::read);
        ArrayHandle<Scalar4> d_pos(m_mpcd_pdata->getPositions(),
                                   access_location::device,
                                   access_mode::read);
        ArrayHandle<Scalar4> d_vel(m_mpcd_pdata->getVelocities(),
                                   access_location::device,
                                   access_mode::read);
        ArrayHandle<unsigned int> d_tag(m_mpcd_pdata->getTags(),
                                        access_location::device,
                                        access_mode::read);
        ArrayHandle<Scalar4> d_pos_alt(m_mpcd_pdata->getAltPositions(),
                                       access_location::device,
                                       access_mode::overwrite);
        ArrayHandle<Scalar4> d_vel_alt(m_mpcd_pdata->getAltVelocities(),
                                       access_location::device,
                                       access_mode::overwrite);
        ArrayHandle<unsigned int> d_tag_alt(m_mpcd_pdata->getAltTags(),
                                            access_location::device,
                                            access_mode::overwrite);
        m_apply_tuner->begin();
        mpcd::gpu::sort_apply(d_pos_alt.data,
                              d_vel_alt.data,
                              d_tag_alt.data,
                              d_pos.data,
                              d_vel.data,
                              d_tag.data,
                              d_order.data,
                              m_mpcd_pdata->getN(),
                              m_apply_tuner->getParam());
        if (m_exec_conf->isCUDAErrorCheckingEnabled())
            CHECK_CUDA_ERROR();
        m_apply_tuner->end();
        if (m_mpcd_pdata->getNVirtual() > 0)
            {
            const unsigned int N = m_mpcd_pdata->getN();
            const unsigned int Nvirtual = m_mpcd_pdata->getNVirtual();
            cudaMemcpyAsync(d_pos_alt.data + N,
                            d_pos.data + N,
                            Nvirtual * sizeof(Scalar4),
                            cudaMemcpyDeviceToDevice);
            cudaMemcpyAsync(d_vel_alt.data + N,
                            d_vel.data + N,
                            Nvirtual * sizeof(Scalar4),
                            cudaMemcpyDeviceToDevice);
            cudaMemcpyAsync(d_tag_alt.data + N,
                            d_tag.data + N,
                            Nvirtual * sizeof(unsigned int),
                            cudaMemcpyDeviceToDevice);
            cudaDeviceSynchronize();
            }
        }
    m_mpcd_pdata->swapPositions();
    m_mpcd_pdata->swapVelocities();
    m_mpcd_pdata->swapTags();
    }
void mpcd :: SorterGPU :: applyOrder ( ) const { { ArrayHandle < unsigned int > d_order ( m_order , access_location :: device , access_mode :: read ) ; ArrayHandle < Scalar4 > d_pos ( m_mpcd_pdata -> getPositions ( ) , access_location :: device , access_mode :: read ) ; ArrayHandle < Scalar4 > d_vel ( m_mpcd_pdata -> getVelocities ( ) , access_location :: device , access_mode :: read ) ; ArrayHandle < unsigned int > d_tag ( m_mpcd_pdata -> getTags ( ) , access_location :: device , access_mode :: read ) ; ArrayHandle < Scalar4 > d_pos_alt ( m_mpcd_pdata -> getAltPositions ( ) , access_location :: device , access_mode :: overwrite ) ; ArrayHandle < Scalar4 > d_vel_alt ( m_mpcd_pdata -> getAltVelocities ( ) , access_location :: device , access_mode :: overwrite ) ; ArrayHandle < unsigned int > d_tag_alt ( m_mpcd_pdata -> getAltTags ( ) , access_location :: device , access_mode :: overwrite ) ; m_apply_tuner -> begin ( ) ; mpcd :: gpu :: sort_apply ( d_pos_alt . data , d_vel_alt . data , d_tag_alt . data , d_pos . data , d_vel . data , d_tag . data , d_order . data , m_mpcd_pdata -> getN ( ) , m_apply_tuner -> getParam ( ) ) ; if ( m_exec_conf -> isCUDAErrorCheckingEnabled ( ) ) CHECK_CUDA_ERROR ( ) ; m_apply_tuner -> end ( ) ; if ( m_mpcd_pdata -> getNVirtual ( ) > 0 ) { const unsigned int N = m_mpcd_pdata -> getN ( ) ; const unsigned int Nvirtual = m_mpcd_pdata -> getNVirtual ( ) ; cudaMemcpyAsync ( d_pos_alt . data + N , d_pos . data + N , Nvirtual * sizeof ( Scalar4 ) , cudaMemcpyDeviceToDevice ) ; cudaMemcpyAsync ( d_vel_alt . data + N , d_vel . data + N , Nvirtual * sizeof ( Scalar4 ) , cudaMemcpyDeviceToDevice ) ; cudaMemcpyAsync ( d_tag_alt . data + N , d_tag . data + N , Nvirtual * sizeof ( unsigned int ) , cudaMemcpyDeviceToDevice ) ; cudaDeviceSynchronize ( ) ; } } m_mpcd_pdata -> swapPositions ( ) ; m_mpcd_pdata -> swapVelocities ( ) ; m_mpcd_pdata -> swapTags ( ) ; }
The sorted order is applied by swapping out the alternate per-particle data
 arrays.
The sorted order is applied by swapping out the alternate per - particle data arrays .
// apply the sorted order // copy virtual particle data if it exists // swap out sorted data