Pin important buffers in CUDA bondeds

This allows for asynchronous transfers.

Change-Id: I0f251c17bdd44856bf773641633e8fa78e6be894

diff --git a/src/gromacs/listed-forces/bonded.cu b/src/gromacs/listed-forces/bonded.cu
index 52cfd70e60abaeaa3c7133fac25201674c8fcbf6..7e6c39e531779b1817b4e527b5c9886b06d47f3e 100644 (file)
--- a/src/gromacs/listed-forces/bonded.cu
+++ b/src/gromacs/listed-forces/bonded.cu
@@ -1037,7 +1037,7 @@ update_gpu_bonded(GpuBondedLists *gpuBondedLists)
     // forceatoms
     for (int ftype : ftypesOnGpu)
     {
-        const InteractionList &iList = gpuBondedLists->iLists[ftype];
+        const auto &iList = gpuBondedLists->iLists[ftype];
 
         if (iList.size() > 0)
         {
@@ -1047,7 +1047,7 @@ update_gpu_bonded(GpuBondedLists *gpuBondedLists)
 
             copyToDeviceBuffer(&iListDevice.iatoms, iList.iatoms.data(),
                                0, iList.size(),
-                               *stream, GpuApiCallBehavior::Sync, nullptr);
+                               *stream, GpuApiCallBehavior::Async, nullptr);
         }
     }
 }
@@ -1082,7 +1082,7 @@ launch_bonded_kernels(t_forcerec   *fr,
 
     for (int ftype : ftypesOnGpu)
     {
-        const InteractionList &iList = gpuBondedLists->iLists[ftype];
+        const auto &iList = gpuBondedLists->iLists[ftype];
 
         if (iList.size() > 0)
         {
@@ -1116,7 +1116,7 @@ launch_bonded_kernels(t_forcerec   *fr,
 
     for (int ftype : ftypesOnGpu)
     {
-        const InteractionList &iList = fr->gpuBondedLists->iLists[ftype];
+        const auto &iList = fr->gpuBondedLists->iLists[ftype];
 
         if (iList.size() > 0)
         {
@@ -1250,7 +1250,7 @@ bonded_gpu_get_energies(t_forcerec *fr,  gmx_enerdata_t *enerd)
     float        *vtot   = gpuBondedLists->vtot.data();
     copyFromDeviceBuffer(vtot, &gpuBondedLists->vtotDevice,
                          0, F_NRE,
-                         *stream, GpuApiCallBehavior::Sync, nullptr);
+                         *stream, GpuApiCallBehavior::Async, nullptr);
     cudaError_t stat = cudaStreamSynchronize(*stream);
     CU_RET_ERR(stat, "D2H transfer failed");
 

diff --git a/src/gromacs/listed-forces/manage-threading.cpp b/src/gromacs/listed-forces/manage-threading.cpp
index a3941bf414573878404e935e5ebaee37f525eb80..ce7a77f6c248ff8c5201af3fd1645ade737886ac 100644 (file)
--- a/src/gromacs/listed-forces/manage-threading.cpp
+++ b/src/gromacs/listed-forces/manage-threading.cpp
@@ -339,7 +339,7 @@ static void divide_bondeds_over_threads(bonded_threading_t *bt,
 }
 //! Converts \p src with atom indices in state order to \p dest in nbnxn order
 static void convertIlistToNbnxnOrder(const t_ilist            &src,
-                                     InteractionList          *dest,
+                                     HostInteractionList      *dest,
                                      int                       numAtomsPerInteraction,
                                      gmx::ArrayRef<const int>  nbnxnAtomOrder)
 {

diff --git a/src/gromacs/listed-forces/manage-threading.h b/src/gromacs/listed-forces/manage-threading.h
index d052045e46ce5b47fd94c15fcaa3ecca171dadfb..051f362ac00e947382eed29e9ab0bb68fbaa1916 100644 (file)
--- a/src/gromacs/listed-forces/manage-threading.h
+++ b/src/gromacs/listed-forces/manage-threading.h
@@ -50,6 +50,7 @@
 
 #include <string>
 
+#include "gromacs/gpu_utils/hostallocator.h"
 #include "gromacs/topology/idef.h"
 #include "gromacs/utility/arrayref.h"
 
@@ -57,14 +58,17 @@ struct bonded_threading_t;
 struct gmx_mtop_t;
 struct t_inputrec;
 
-/*! \brief List of all bonded function types supported on a GPUs
+/*! \brief The number on bonded function types supported on GPUs */
+constexpr int c_numFtypesOnGpu = 8;
+
+/*! \brief List of all bonded function types supported on GPUs
  *
  * \note This list should be in sync with the actual GPU code.
  * \note Perturbed interactions are not supported on GPUs.
  * \note The function types in the list are ordered on increasing value.
  * \note Currently bonded are only supported with CUDA, not with OpenCL.
  */
-constexpr std::array<int, 8> ftypesOnGpu =
+constexpr std::array<int, c_numFtypesOnGpu> ftypesOnGpu =
 {
     F_BONDS,
     F_ANGLES,
@@ -76,6 +80,22 @@ constexpr std::array<int, 8> ftypesOnGpu =
     F_LJ14
 };
 
+/*! \libinternal \brief Version of InteractionList that supports pinning */
+struct HostInteractionList
+{
+    /*! \brief Returns the total number of elements in iatoms */
+    int size() const
+    {
+        return iatoms.size();
+    }
+
+    /*! \brief List of interactions, see explanation further down */
+    std::vector < int, gmx::HostAllocator < int>> iatoms = {{}, gmx::HostAllocationPolicy(gmx::PinningPolicy::PinnedIfSupported)};
+};
+
+/*! \brief Convenience alias for set of pinned interaction lists */
+using HostInteractionLists = std::array<HostInteractionList, F_NRE>;
+
 /*! \internal \brief Struct for storing lists of bonded interaction for evaluation on a GPU */
 struct GpuBondedLists
 {
@@ -98,15 +118,19 @@ struct GpuBondedLists
     }
 #endif
 
-    InteractionLists    iLists;                      /**< The interaction lists */
-    bool                haveInteractions;            /**< Tells whether there are any interaction in iLists */
+    HostInteractionLists  iLists;                      /**< The interaction lists */
+    bool                  haveInteractions;            /**< Tells whether there are any interaction in iLists */
+
+    t_iparams            *forceparamsDevice = nullptr; /**< Bonded parameters for device-side use */
+    t_ilist               iListsDevice[F_NRE];         /**< Interaction lists on the device */
 
-    t_iparams          *forceparamsDevice = nullptr; /**< Bonded parameters for device-side use */
-    t_ilist             iListsDevice[F_NRE];         /**< Interaction lists on the device */
-    std::vector<float>  vtot;                        /**< Host-side virial buffer */
-    float              *vtotDevice   = nullptr;      /**< Device-side total virial */
+    //! \brief Host-side virial buffer
+    std::vector < float, gmx::HostAllocator < float>> vtot = {{}, gmx::HostAllocationPolicy(gmx::PinningPolicy::PinnedIfSupported)};
+    //! \brief Device-side total virial
+    float                *vtotDevice   = nullptr;
 
-    void               *stream;                      /**< Bonded GPU stream */
+    //! \brief Bonded GPU stream
+    void                 *stream;
 };