src/gromacs/nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2018,2019, 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  *
  38  * \brief
  39  * CUDA kernels for GPU versions of copy_rvec_to_nbat_real and add_nbat_f_to_f.
  40  *
  41  *  \author Alan Gray <alang@nvidia.com>
  42  *  \author Jon Vincent <jvincent@nvidia.com>
  43  */
  44
  45 #include "gromacs/gpu_utils/vectype_ops.cuh"
  46 #include "gromacs/nbnxm/nbnxm.h"
  47
  48 /*! \brief CUDA kernel for transforming position coordinates from rvec to nbnxm layout.
  49  *
  50  * TODO:
  51  *  - improve/simplify/document use of cxy_na and na_round
  52  *  - rename kernel so naming matches with the other NBNXM kernels;
  53  *  - enable separate compilation unit
  54
  55  * \param[in]     numColumns        extent of cell-level parallelism
  56  * \param[out]    xnb               position buffer in nbnxm layout
  57  * \param[in]     setFillerCoords   tells whether to set the coordinates of the filler particles
  58  * \param[in]     x                 position buffer
  59  * \param[in]     a                 atom index mapping stride between atoms in memory
  60  * \param[in]     cxy_na            array of extents
  61  * \param[in]     cxy_ind           array of cell indices
  62  * \param[in]     cellOffset        first cell
  63  * \param[in]     numAtomsPerCell   number of atoms per cell
  64  */
  65 __global__ void nbnxn_gpu_x_to_nbat_x_kernel(int numColumns,
  66                                              float* __restrict__ xnb,
  67                                              bool setFillerCoords,
  68                                              const rvec* __restrict__ x,
  69                                              const int* __restrict__ a,
  70                                              const int* __restrict__ cxy_na,
  71                                              const int* __restrict__ cxy_ind,
  72                                              int cellOffset,
  73                                              int numAtomsPerCell);
  74
  75
  76 __global__ void nbnxn_gpu_x_to_nbat_x_kernel(int numColumns,
  77                                              float* __restrict__ xnb,
  78                                              bool setFillerCoords,
  79                                              const rvec* __restrict__ x,
  80                                              const int* __restrict__ a,
  81                                              const int* __restrict__ cxy_na,
  82                                              const int* __restrict__ cxy_ind,
  83                                              int cellOffset,
  84                                              int numAtomsPerCell)
  85 {
  86
  87
  88     const float farAway = -1000000.0f;
  89
  90     /* map cell-level parallelism to y component of CUDA block index */
  91     int cxy = blockIdx.y;
  92
  93     if (cxy < numColumns)
  94     {
  95
  96         int na = cxy_na[cxy];
  97         int a0 = (cellOffset + cxy_ind[cxy]) * numAtomsPerCell;
  98         int na_round;
  99         if (setFillerCoords)
 100         {
 101             // TODO: This can be done more efficiently
 102             na_round = (cxy_ind[cxy + 1] - cxy_ind[cxy]) * numAtomsPerCell;
 103         }
 104         else
 105         {
 106             /* We fill only the real particle locations.
 107              * We assume the filling entries at the end have been
 108              * properly set before during pair-list generation.
 109              */
 110             na_round = na;
 111         }
 112
 113         /* map parallelism within a cell to x component of CUDA block index linearized
 114          * with threads within a block */
 115         int i, j0;
 116         i = blockIdx.x * blockDim.x + threadIdx.x;
 117
 118         j0 = a0 * STRIDE_XYZQ;
 119
 120         // destination address where x shoud be stored in nbnxm layout
 121         float3* x_dest = (float3*)&xnb[j0 + 4 * i];
 122
 123         /* perform conversion of each element */
 124         if (i < na_round)
 125         {
 126             if (i < na)
 127             {
 128                 *x_dest = *((float3*)x[a[a0 + i]]);
 129             }
 130             else
 131             {
 132                 *x_dest = make_float3(farAway);
 133             }
 134         }
 135     }
 136 }
 137
 138 /*! \brief CUDA kernel to sum up the force components
 139  *
 140  * \tparam        accumulateForce  If the initial forces in \p d_fTotal should be saved.
 141  * \tparam        addPmeForce      Whether the PME force should be added to the total.
 142  *
 143  * \param[in]     d_fNB            Non-bonded forces in nbat format.
 144  * \param[in]     d_fPme           PME forces.
 145  * \param[in,out] d_fTotal         Force buffer to be reduced into.
 146  * \param[in]     cell             Cell index mapping.
 147  * \param[in]     atomStart        Start atom index.
 148  * \param[in]     numAtoms         Number of atoms.
 149  */
 150 template<bool accumulateForce, bool addPmeForce>
 151 __global__ void nbnxn_gpu_add_nbat_f_to_f_kernel(const float3* __restrict__ d_fNB,
 152                                                  const float3* __restrict__ d_fPme,
 153                                                  float3* d_fTotal,
 154                                                  const int* __restrict__ d_cell,
 155                                                  const int atomStart,
 156                                                  const int numAtoms);
 157 template<bool accumulateForce, bool addPmeForce>
 158 __global__ void nbnxn_gpu_add_nbat_f_to_f_kernel(const float3* __restrict__ d_fNB,
 159                                                  const float3* __restrict__ d_fPme,
 160                                                  float3* d_fTotal,
 161                                                  const int* __restrict__ d_cell,
 162                                                  const int atomStart,
 163                                                  const int numAtoms)
 164 {
 165
 166     /* map particle-level parallelism to 1D CUDA thread and block index */
 167     int threadIndex = blockIdx.x * blockDim.x + threadIdx.x;
 168
 169     /* perform addition for each particle*/
 170     if (threadIndex < numAtoms)
 171     {
 172
 173         int     i     = d_cell[atomStart + threadIndex];
 174         float3* fDest = (float3*)&d_fTotal[atomStart + threadIndex];
 175         float3  temp;
 176
 177         if (accumulateForce)
 178         {
 179             temp = *fDest;
 180             temp += d_fNB[i];
 181         }
 182         else
 183         {
 184             temp = d_fNB[i];
 185         }
 186         if (addPmeForce)
 187         {
 188             temp += d_fPme[atomStart + threadIndex];
 189         }
 190         *fDest = temp;
 191     }
 192     return;
 193 }