src/gromacs/ewald/pme_coordinate_receiver_gpu_impl.cu

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2019,2020,2021, by the GROMACS development team, led by
   5  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   6  * and including many others, as listed in the AUTHORS file in the
   7  * top-level source directory and at http://www.gromacs.org.
   8  *
   9  * GROMACS is free software; you can redistribute it and/or
  10  * modify it under the terms of the GNU Lesser General Public License
  11  * as published by the Free Software Foundation; either version 2.1
  12  * of the License, or (at your option) any later version.
  13  *
  14  * GROMACS is distributed in the hope that it will be useful,
  15  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  16  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  17  * Lesser General Public License for more details.
  18  *
  19  * You should have received a copy of the GNU Lesser General Public
  20  * License along with GROMACS; if not, see
  21  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  22  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  23  *
  24  * If you want to redistribute modifications to GROMACS, please
  25  * consider that scientific software is very special. Version
  26  * control is crucial - bugs must be traceable. We will be happy to
  27  * consider code for inclusion in the official distribution, but
  28  * derived work must not be called official GROMACS. Details are found
  29  * in the README & COPYING files - if they are missing, get the
  30  * official version at http://www.gromacs.org.
  31  *
  32  * To help us fund GROMACS development, we humbly ask that you cite
  33  * the research papers on the package. Check out http://www.gromacs.org.
  34  */
  35 /*! \internal \file
  36  *
  37  * \brief Implements class which recieves coordinates to GPU memory on PME task using CUDA
  38  *
  39  *
  40  * \author Alan Gray <alang@nvidia.com>
  41  *
  42  * \ingroup module_ewald
  43  */
  44 #include "gmxpre.h"
  45
  46 #include "gromacs/ewald/pme_pp_communication.h"
  47 #include "pme_coordinate_receiver_gpu_impl.h"
  48
  49 #include "config.h"
  50
  51 #include "gromacs/ewald/pme_force_sender_gpu.h"
  52 #include "gromacs/gpu_utils/cudautils.cuh"
  53 #include "gromacs/gpu_utils/gpueventsynchronizer.h"
  54 #include "gromacs/utility/gmxmpi.h"
  55
  56 namespace gmx
  57 {
  58
  59 PmeCoordinateReceiverGpu::Impl::Impl(MPI_Comm                     comm,
  60                                      const DeviceContext&         deviceContext,
  61                                      gmx::ArrayRef<const PpRanks> ppRanks) :
  62     comm_(comm), requests_(ppRanks.size(), MPI_REQUEST_NULL), deviceContext_(deviceContext)
  63 {
  64     // Create streams to manage pipelining
  65     ppCommManagers_.reserve(ppRanks.size());
  66     for (auto& ppRank : ppRanks)
  67     {
  68         ppCommManagers_.emplace_back(PpCommManager{
  69                 ppRank,
  70                 std::make_unique<DeviceStream>(deviceContext_, DeviceStreamPriority::High, false),
  71                 nullptr,
  72                 { 0, 0 } });
  73     }
  74 }
  75
  76 PmeCoordinateReceiverGpu::Impl::~Impl() = default;
  77
  78 void PmeCoordinateReceiverGpu::Impl::reinitCoordinateReceiver(DeviceBuffer<RVec> d_x)
  79 {
  80     int indEnd = 0;
  81     for (auto& ppCommManager : ppCommManagers_)
  82     {
  83         int indStart = indEnd;
  84         indEnd       = indStart + ppCommManager.ppRank.numAtoms;
  85
  86         ppCommManager.atomRange = std::make_tuple(indStart, indEnd);
  87
  88         // Need to send address to PP rank only for thread-MPI as PP rank pushes data using cudamemcpy
  89         if (GMX_THREAD_MPI)
  90         {
  91             // Data will be transferred directly from GPU.
  92             void* sendBuf = reinterpret_cast<void*>(&d_x[indStart]);
  93 #if GMX_MPI
  94             MPI_Send(&sendBuf, sizeof(void**), MPI_BYTE, ppCommManager.ppRank.rankId, 0, comm_);
  95 #else
  96             GMX_UNUSED_VALUE(sendBuf);
  97 #endif
  98         }
  99     }
 100 }
 101
 102 /*! \brief Receive coordinate synchronizer pointer from the PP ranks. */
 103 void PmeCoordinateReceiverGpu::Impl::receiveCoordinatesSynchronizerFromPpCudaDirect(int ppRank)
 104 {
 105     GMX_ASSERT(GMX_THREAD_MPI,
 106                "receiveCoordinatesSynchronizerFromPpCudaDirect is expected to be called only for "
 107                "Thread-MPI");
 108
 109     // Data will be pushed directly from PP task
 110
 111 #if GMX_MPI
 112     // Receive event from PP task
 113     // NOLINTNEXTLINE(bugprone-sizeof-expression)
 114     MPI_Irecv(&ppCommManagers_[ppRank].sync,
 115               sizeof(GpuEventSynchronizer*),
 116               MPI_BYTE,
 117               ppRank,
 118               0,
 119               comm_,
 120               &(requests_[ppRank]));
 121 #else
 122     GMX_UNUSED_VALUE(ppRank);
 123 #endif
 124 }
 125
 126 /*! \brief Receive coordinate data using CUDA-aware MPI */
 127 void PmeCoordinateReceiverGpu::Impl::launchReceiveCoordinatesFromPpCudaMpi(DeviceBuffer<RVec> recvbuf,
 128                                                                            int numAtoms,
 129                                                                            int numBytes,
 130                                                                            int ppRank)
 131 {
 132     GMX_ASSERT(GMX_LIB_MPI,
 133                "launchReceiveCoordinatesFromPpCudaMpi is expected to be called only for Lib-MPI");
 134
 135 #if GMX_MPI
 136     MPI_Irecv(&recvbuf[numAtoms], numBytes, MPI_BYTE, ppRank, eCommType_COORD_GPU, comm_, &(requests_[ppRank]));
 137 #else
 138     GMX_UNUSED_VALUE(recvbuf);
 139     GMX_UNUSED_VALUE(numAtoms);
 140     GMX_UNUSED_VALUE(numBytes);
 141     GMX_UNUSED_VALUE(ppRank);
 142 #endif
 143 }
 144
 145 int PmeCoordinateReceiverGpu::Impl::synchronizeOnCoordinatesFromPpRank(int pipelineStage,
 146                                                                        const DeviceStream& deviceStream)
 147 {
 148 #if GMX_MPI
 149     int senderRank = -1; // Rank of PP task that is associated with this invocation.
 150 #    if (!GMX_THREAD_MPI)
 151     // Wait on data from any one of the PP sender GPUs
 152     MPI_Waitany(requests_.size(), requests_.data(), &senderRank, MPI_STATUS_IGNORE);
 153     GMX_ASSERT(senderRank >= 0, "Rank of sending PP task must be 0 or greater");
 154     GMX_UNUSED_VALUE(pipelineStage);
 155     GMX_UNUSED_VALUE(deviceStream);
 156 #    else
 157     // MPI_Waitany is not available in thread-MPI. However, the
 158     // MPI_Wait here is not associated with data but is host-side
 159     // scheduling code to receive a CUDA event, and will be executed
 160     // in advance of the actual data transfer. Therefore we can
 161     // receive in order of pipeline stage, still allowing the
 162     // scheduled GPU-direct comms to initiate out-of-order in their
 163     // respective streams. For cases with CPU force computations, the
 164     // scheduling is less asynchronous (done on a per-step basis), so
 165     // host-side improvements should be investigated as tracked in
 166     // issue #4047
 167     senderRank = pipelineStage;
 168     MPI_Wait(&(requests_[senderRank]), MPI_STATUS_IGNORE);
 169     ppCommManagers_[senderRank].sync->enqueueWaitEvent(deviceStream);
 170 #    endif
 171     return senderRank;
 172 #endif
 173 }
 174
 175 void PmeCoordinateReceiverGpu::Impl::synchronizeOnCoordinatesFromAllPpRanks(const DeviceStream& deviceStream)
 176 {
 177     for (int i = 0; i < static_cast<int>(ppCommManagers_.size()); i++)
 178     {
 179         synchronizeOnCoordinatesFromPpRank(i, deviceStream);
 180     }
 181 }
 182 DeviceStream* PmeCoordinateReceiverGpu::Impl::ppCommStream(int senderIndex)
 183 {
 184     return ppCommManagers_[senderIndex].stream.get();
 185 }
 186
 187 std::tuple<int, int> PmeCoordinateReceiverGpu::Impl::ppCommAtomRange(int senderIndex)
 188 {
 189     return ppCommManagers_[senderIndex].atomRange;
 190 }
 191
 192 int PmeCoordinateReceiverGpu::Impl::ppCommNumSenderRanks()
 193 {
 194     return ppCommManagers_.size();
 195 }
 196
 197 PmeCoordinateReceiverGpu::PmeCoordinateReceiverGpu(MPI_Comm               comm,
 198                                                    const DeviceContext&   deviceContext,
 199                                                    gmx::ArrayRef<PpRanks> ppRanks) :
 200     impl_(new Impl(comm, deviceContext, ppRanks))
 201 {
 202 }
 203
 204 PmeCoordinateReceiverGpu::~PmeCoordinateReceiverGpu() = default;
 205
 206 void PmeCoordinateReceiverGpu::reinitCoordinateReceiver(DeviceBuffer<RVec> d_x)
 207 {
 208     impl_->reinitCoordinateReceiver(d_x);
 209 }
 210
 211 void PmeCoordinateReceiverGpu::receiveCoordinatesSynchronizerFromPpCudaDirect(int ppRank)
 212 {
 213     impl_->receiveCoordinatesSynchronizerFromPpCudaDirect(ppRank);
 214 }
 215
 216 void PmeCoordinateReceiverGpu::launchReceiveCoordinatesFromPpCudaMpi(DeviceBuffer<RVec> recvbuf,
 217                                                                      int                numAtoms,
 218                                                                      int                numBytes,
 219                                                                      int                ppRank)
 220 {
 221     impl_->launchReceiveCoordinatesFromPpCudaMpi(recvbuf, numAtoms, numBytes, ppRank);
 222 }
 223
 224 int PmeCoordinateReceiverGpu::synchronizeOnCoordinatesFromPpRank(int                 senderIndex,
 225                                                                  const DeviceStream& deviceStream)
 226 {
 227     return impl_->synchronizeOnCoordinatesFromPpRank(senderIndex, deviceStream);
 228 }
 229
 230 void PmeCoordinateReceiverGpu::synchronizeOnCoordinatesFromAllPpRanks(const DeviceStream& deviceStream)
 231 {
 232     impl_->synchronizeOnCoordinatesFromAllPpRanks(deviceStream);
 233 }
 234
 235 DeviceStream* PmeCoordinateReceiverGpu::ppCommStream(int senderIndex)
 236 {
 237     return impl_->ppCommStream(senderIndex);
 238 }
 239
 240 std::tuple<int, int> PmeCoordinateReceiverGpu::ppCommAtomRange(int senderIndex)
 241 {
 242     return impl_->ppCommAtomRange(senderIndex);
 243 }
 244
 245 int PmeCoordinateReceiverGpu::ppCommNumSenderRanks()
 246 {
 247     return impl_->ppCommNumSenderRanks();
 248 }
 249
 250
 251 } // namespace gmx