src/gromacs/nbnxm/sycl/nbnxm_sycl_data_mgmt.cpp

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 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
  36 /*! \internal \file
  37  *  \brief
  38  *  Stubs of functions that must be defined by nbnxm sycl implementation.
  39  *
  40  *  \ingroup module_nbnxm
  41  */
  42 #include "gmxpre.h"
  43
  44 #include "gromacs/gpu_utils/device_stream_manager.h"
  45 #include "gromacs/gpu_utils/pmalloc.h"
  46 #include "gromacs/hardware/device_information.h"
  47 #include "gromacs/mdtypes/interaction_const.h"
  48 #include "gromacs/nbnxm/atomdata.h"
  49 #include "gromacs/nbnxm/gpu_data_mgmt.h"
  50 #include "gromacs/nbnxm/nbnxm_gpu.h"
  51 #include "gromacs/nbnxm/nbnxm_gpu_data_mgmt.h"
  52 #include "gromacs/pbcutil/ishift.h"
  53 #include "gromacs/utility/exceptions.h"
  54
  55 #include "nbnxm_sycl_types.h"
  56
  57 namespace Nbnxm
  58 {
  59
  60 //! This function is documented in the header file
  61 void gpu_clear_outputs(NbnxmGpu* nb, bool computeVirial)
  62 {
  63     NBAtomData*         adat        = nb->atdat;
  64     const DeviceStream& localStream = *nb->deviceStreams[InteractionLocality::Local];
  65     // Clear forces
  66     clearDeviceBufferAsync(&adat->f, 0, nb->atdat->numAtoms, localStream);
  67     // Clear shift force array and energies if the outputs were used in the current step
  68     if (computeVirial)
  69     {
  70         clearDeviceBufferAsync(&adat->fShift, 0, SHIFTS, localStream);
  71         clearDeviceBufferAsync(&adat->eLJ, 0, 1, localStream);
  72         clearDeviceBufferAsync(&adat->eElec, 0, 1, localStream);
  73     }
  74 }
  75
  76 /*! \brief Initialize \p atomdata first time; it only gets filled at pair-search. */
  77 static void initAtomdataFirst(NbnxmGpu* nb, int numTypes, const DeviceContext& deviceContext)
  78 {
  79     const DeviceStream& localStream = *nb->deviceStreams[InteractionLocality::Local];
  80     NBAtomData*         atomdata    = nb->atdat;
  81     atomdata->numTypes              = numTypes;
  82     allocateDeviceBuffer(&atomdata->shiftVec, SHIFTS, deviceContext);
  83     atomdata->shiftVecUploaded = false;
  84
  85     allocateDeviceBuffer(&atomdata->fShift, SHIFTS, deviceContext);
  86     allocateDeviceBuffer(&atomdata->eLJ, 1, deviceContext);
  87     allocateDeviceBuffer(&atomdata->eElec, 1, deviceContext);
  88
  89     clearDeviceBufferAsync(&atomdata->fShift, 0, SHIFTS, localStream);
  90     clearDeviceBufferAsync(&atomdata->eElec, 0, 1, localStream);
  91     clearDeviceBufferAsync(&atomdata->eLJ, 0, 1, localStream);
  92
  93     /* initialize to nullptr pointers to data that is not allocated here and will
  94        need reallocation in later */
  95     atomdata->xq = nullptr;
  96     atomdata->f  = nullptr;
  97
  98     /* size -1 indicates that the respective array hasn't been initialized yet */
  99     atomdata->numAtoms      = -1;
 100     atomdata->numAtomsAlloc = -1;
 101 }
 102
 103 /*! \brief Initialize the nonbonded parameter data structure. */
 104 static void initNbparam(NBParamGpu*                     nbp,
 105                         const interaction_const_t&      ic,
 106                         const PairlistParams&           listParams,
 107                         const nbnxn_atomdata_t::Params& nbatParams,
 108                         const DeviceContext&            deviceContext)
 109 {
 110     const int numTypes = nbatParams.numTypes;
 111
 112     set_cutoff_parameters(nbp, &ic, listParams);
 113
 114     nbp->vdwType  = nbnxmGpuPickVdwKernelType(&ic, nbatParams.ljCombinationRule);
 115     nbp->elecType = nbnxmGpuPickElectrostaticsKernelType(&ic, deviceContext.deviceInfo());
 116
 117     /* generate table for PME */
 118     nbp->coulomb_tab = nullptr;
 119     if (nbp->elecType == ElecType::EwaldTab || nbp->elecType == ElecType::EwaldTabTwin)
 120     {
 121         GMX_RELEASE_ASSERT(ic.coulombEwaldTables, "Need valid Coulomb Ewald correction tables");
 122         init_ewald_coulomb_force_table(*ic.coulombEwaldTables, nbp, deviceContext);
 123     }
 124
 125     /* set up LJ parameter lookup table */
 126     if (!useLjCombRule(nbp->vdwType))
 127     {
 128         static_assert(sizeof(decltype(nbp->nbfp)) == 2 * sizeof(decltype(*nbatParams.nbfp.data())),
 129                       "Mismatch in the size of host / device data types");
 130         initParamLookupTable(&nbp->nbfp,
 131                              &nbp->nbfp_texobj,
 132                              reinterpret_cast<const Float2*>(nbatParams.nbfp.data()),
 133                              numTypes * numTypes,
 134                              deviceContext);
 135     }
 136
 137     /* set up LJ-PME parameter lookup table */
 138     if (ic.vdwtype == VanDerWaalsType::Pme)
 139     {
 140         static_assert(sizeof(decltype(nbp->nbfp_comb))
 141                               == 2 * sizeof(decltype(*nbatParams.nbfp_comb.data())),
 142                       "Mismatch in the size of host / device data types");
 143         initParamLookupTable(&nbp->nbfp_comb,
 144                              &nbp->nbfp_comb_texobj,
 145                              reinterpret_cast<const Float2*>(nbatParams.nbfp_comb.data()),
 146                              numTypes,
 147                              deviceContext);
 148     }
 149 }
 150
 151 NbnxmGpu* gpu_init(const gmx::DeviceStreamManager& deviceStreamManager,
 152                    const interaction_const_t*      ic,
 153                    const PairlistParams&           listParams,
 154                    const nbnxn_atomdata_t*         nbat,
 155                    const bool                      bLocalAndNonlocal)
 156 {
 157     auto* nb                              = new NbnxmGpu();
 158     nb->deviceContext_                    = &deviceStreamManager.context();
 159     nb->atdat                             = new NBAtomData;
 160     nb->nbparam                           = new NBParamGpu;
 161     nb->plist[InteractionLocality::Local] = new Nbnxm::gpu_plist;
 162     if (bLocalAndNonlocal)
 163     {
 164         nb->plist[InteractionLocality::NonLocal] = new Nbnxm::gpu_plist;
 165     }
 166
 167     nb->bUseTwoStreams = bLocalAndNonlocal;
 168
 169     nb->timers  = nullptr;
 170     nb->timings = nullptr;
 171
 172     /* init nbst */
 173     pmalloc(reinterpret_cast<void**>(&nb->nbst.eLJ), sizeof(*nb->nbst.eLJ));
 174     pmalloc(reinterpret_cast<void**>(&nb->nbst.eElec), sizeof(*nb->nbst.eElec));
 175     pmalloc(reinterpret_cast<void**>(&nb->nbst.fShift), SHIFTS * sizeof(*nb->nbst.fShift));
 176
 177     init_plist(nb->plist[InteractionLocality::Local]);
 178
 179     /* local/non-local GPU streams */
 180     GMX_RELEASE_ASSERT(deviceStreamManager.streamIsValid(gmx::DeviceStreamType::NonBondedLocal),
 181                        "Local non-bonded stream should be initialized to use GPU for non-bonded.");
 182     nb->deviceStreams[InteractionLocality::Local] =
 183             &deviceStreamManager.stream(gmx::DeviceStreamType::NonBondedLocal);
 184     // In general, it's not strictly necessary to use 2 streams for SYCL, since they are
 185     // out-of-order. But for the time being, it will be less disruptive to keep them.
 186     if (nb->bUseTwoStreams)
 187     {
 188         init_plist(nb->plist[InteractionLocality::NonLocal]);
 189
 190         GMX_RELEASE_ASSERT(deviceStreamManager.streamIsValid(gmx::DeviceStreamType::NonBondedNonLocal),
 191                            "Non-local non-bonded stream should be initialized to use GPU for "
 192                            "non-bonded with domain decomposition.");
 193         nb->deviceStreams[InteractionLocality::NonLocal] =
 194                 &deviceStreamManager.stream(gmx::DeviceStreamType::NonBondedNonLocal);
 195     }
 196
 197     nb->bDoTime = false;
 198
 199     const nbnxn_atomdata_t::Params& nbatParams    = nbat->params();
 200     const DeviceContext&            deviceContext = *nb->deviceContext_;
 201
 202     initNbparam(nb->nbparam, *ic, listParams, nbatParams, deviceContext);
 203     initAtomdataFirst(nb, nbatParams.numTypes, deviceContext);
 204
 205     return nb;
 206 }
 207
 208 void gpu_upload_shiftvec(NbnxmGpu* nb, const nbnxn_atomdata_t* nbatom)
 209 {
 210     NBAtomData*         adat        = nb->atdat;
 211     const DeviceStream& localStream = *nb->deviceStreams[InteractionLocality::Local];
 212
 213     /* only if we have a dynamic box */
 214     if (nbatom->bDynamicBox || !adat->shiftVecUploaded)
 215     {
 216         GMX_ASSERT(adat->shiftVec.elementSize() == sizeof(nbatom->shift_vec[0]),
 217                    "Sizes of host- and device-side shift vectors should be the same.");
 218         copyToDeviceBuffer(&adat->shiftVec,
 219                            reinterpret_cast<const Float3*>(nbatom->shift_vec.data()),
 220                            0,
 221                            SHIFTS,
 222                            localStream,
 223                            GpuApiCallBehavior::Async,
 224                            nullptr);
 225         adat->shiftVecUploaded = true;
 226     }
 227 }
 228
 229 void gpu_free(NbnxmGpu* nb)
 230 {
 231     if (nb == nullptr)
 232     {
 233         return;
 234     }
 235
 236     NBAtomData* atdat   = nb->atdat;
 237     NBParamGpu* nbparam = nb->nbparam;
 238
 239     if ((!nbparam->coulomb_tab)
 240         && (nbparam->elecType == ElecType::EwaldTab || nbparam->elecType == ElecType::EwaldTabTwin))
 241     {
 242         destroyParamLookupTable(&nbparam->coulomb_tab, nbparam->coulomb_tab_texobj);
 243     }
 244
 245     if (!useLjCombRule(nb->nbparam->vdwType))
 246     {
 247         destroyParamLookupTable(&nbparam->nbfp, nbparam->nbfp_texobj);
 248     }
 249
 250     if (nbparam->vdwType == VdwType::EwaldGeom || nbparam->vdwType == VdwType::EwaldLB)
 251     {
 252         destroyParamLookupTable(&nbparam->nbfp_comb, nbparam->nbfp_comb_texobj);
 253     }
 254
 255     /* Free plist */
 256     auto* plist = nb->plist[InteractionLocality::Local];
 257     delete plist;
 258     if (nb->bUseTwoStreams)
 259     {
 260         auto* plist_nl = nb->plist[InteractionLocality::NonLocal];
 261         delete plist_nl;
 262     }
 263
 264     /* Free nbst */
 265     pfree(nb->nbst.eLJ);
 266     nb->nbst.eLJ = nullptr;
 267
 268     pfree(nb->nbst.eElec);
 269     nb->nbst.eElec = nullptr;
 270
 271     pfree(nb->nbst.fShift);
 272     nb->nbst.fShift = nullptr;
 273
 274     delete atdat;
 275     delete nbparam;
 276     delete nb;
 277 }
 278
 279 int gpu_min_ci_balanced(NbnxmGpu* nb)
 280 {
 281     // SYCL-TODO: Logic and magic values taken from OpenCL
 282     static constexpr unsigned int balancedFactor = 50;
 283     if (nb == nullptr)
 284     {
 285         return 0;
 286     }
 287     const cl::sycl::device device = nb->deviceContext_->deviceInfo().syclDevice;
 288     const int numComputeUnits     = device.get_info<cl::sycl::info::device::max_compute_units>();
 289     return balancedFactor * numComputeUnits;
 290 }
 291
 292 } // namespace Nbnxm