src/gromacs/pbcutil/pbc_aiuc_cuda.cuh

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 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 /*! \libinternal \file
  36  *
  37  * \brief Basic routines to handle periodic boundary conditions with CUDA.
  38  *
  39  * This file contains GPU implementation of the PBC-aware vector evaluation.
  40  *
  41  * \todo CPU, GPU and SIMD routines essentially do the same operations on
  42  *       different data-types. Currently this leads to code duplication,
  43  *       which has to be resolved. For details, see Redmine task #2863
  44  *       https://redmine.gromacs.org/issues/2863
  45  *
  46  * \author Mark Abraham <mark.j.abraham@gmail.com>
  47  * \author Berk Hess <hess@kth.se>
  48  * \author Artem Zhmurov <zhmurov@gmail.com>
  49  *
  50  * \inlibraryapi
  51  * \ingroup module_pbcutil
  52  */
  53 #ifndef GMX_PBCUTIL_PBC_AIUC_CUDA_CUH
  54 #define GMX_PBCUTIL_PBC_AIUC_CUDA_CUH
  55
  56 #include "gromacs/gpu_utils/gpu_vec.cuh"
  57 #include "gromacs/gpu_utils/vectype_ops.cuh"
  58 #include "gromacs/pbcutil/pbc_aiuc.h"
  59
  60 /*! \brief Computes the vector between two points taking PBC into account.
  61  *
  62  * Computes the vector dr between points r2 and r1, taking into account the
  63  * periodic boundary conditions, described in pbcAiuc object. Note that this
  64  * routine always does the PBC arithmetic for all directions, multiplying the
  65  * displacements by zeroes if the corresponding direction is not periodic.
  66  * For triclinic boxes only distances up to half the smallest box diagonal
  67  * element are guaranteed to be the shortest. This means that distances from
  68  * 0.5/sqrt(2) times a box vector length (e.g. for a rhombic dodecahedron)
  69  * can use a more distant periodic image.
  70  *
  71  * \todo This routine uses CUDA float4 types for input coordinates and
  72  *       returns in rvec data-type. Other than that, it does essentially
  73  *       the same thing as the version below, as well as SIMD and CPU
  74  *       versions. This routine is used in gpubonded module.
  75  *       To avoid code duplication, these implementations should be
  76  *       unified. See Redmine task #2863:
  77  *       https://redmine.gromacs.org/issues/2863
  78  *
  79  * \param[in]  pbcAiuc  PBC object.
  80  * \param[in]  r1       Coordinates of the first point.
  81  * \param[in]  r2       Coordinates of the second point.
  82  * \param[out] dr       Resulting distance.
  83  */
  84 template <bool returnShift>
  85 static __forceinline__ __device__
  86 int pbcDxAiuc(const PbcAiuc &pbcAiuc,
  87               const float4  &r1,
  88               const float4  &r2,
  89               fvec           dr)
  90 {
  91     dr[XX] = r1.x - r2.x;
  92     dr[YY] = r1.y - r2.y;
  93     dr[ZZ] = r1.z - r2.z;
  94
  95     float shz  = rintf(dr[ZZ]*pbcAiuc.invBoxDiagZ);
  96     dr[XX]    -= shz*pbcAiuc.boxZX;
  97     dr[YY]    -= shz*pbcAiuc.boxZY;
  98     dr[ZZ]    -= shz*pbcAiuc.boxZZ;
  99
 100     float shy  = rintf(dr[YY]*pbcAiuc.invBoxDiagY);
 101     dr[XX]    -= shy*pbcAiuc.boxYX;
 102     dr[YY]    -= shy*pbcAiuc.boxYY;
 103
 104     float shx  = rintf(dr[XX]*pbcAiuc.invBoxDiagX);
 105     dr[XX]    -= shx*pbcAiuc.boxXX;
 106
 107     if (returnShift)
 108     {
 109         ivec ishift;
 110
 111         ishift[XX] = -__float2int_rn(shx);
 112         ishift[YY] = -__float2int_rn(shy);
 113         ishift[ZZ] = -__float2int_rn(shz);
 114
 115         return IVEC2IS(ishift);
 116     }
 117     else
 118     {
 119         return 0;
 120     }
 121 }
 122
 123 /*! \brief Computes the vector between two points taking PBC into account.
 124  *
 125  * Computes the vector dr between points r2 and r1, taking into account the
 126  * periodic boundary conditions, described in pbcAiuc object. Same as above,
 127  * only takes and returns data in float3 format. Does not return shifts.
 128  *
 129  * \todo This routine uses CUDA float3 types for both input and returns
 130  *       values. Other than that, it does essentially the same thing as the
 131  *       version above, as well as SIMD and CPU versions. This routine is
 132  *       used in GPU-based constraints.
 133  *       To avoid code duplication, these implementations should be
 134  *       unified. See Redmine task #2863:
 135  *       https://redmine.gromacs.org/issues/2863
 136  *
 137  * \param[in]  pbcAiuc  PBC object.
 138  * \param[in]  r1       Coordinates of the first point.
 139  * \param[in]  r2       Coordinates of the second point.
 140  * \returns    dr       Resulting distance.
 141  */
 142 static __forceinline__ __host__ __device__
 143 float3 pbcDxAiuc(const PbcAiuc &pbcAiuc,
 144                  const float3  &r1,
 145                  const float3  &r2)
 146 {
 147     float3 dr = r1 - r2;
 148
 149     float  shz  = rintf(dr.z*pbcAiuc.invBoxDiagZ);
 150     dr.x    -= shz*pbcAiuc.boxZX;
 151     dr.y    -= shz*pbcAiuc.boxZY;
 152     dr.z    -= shz*pbcAiuc.boxZZ;
 153
 154     float shy  = rintf(dr.y*pbcAiuc.invBoxDiagY);
 155     dr.x    -= shy*pbcAiuc.boxYX;
 156     dr.y    -= shy*pbcAiuc.boxYY;
 157
 158     float shx  = rintf(dr.x*pbcAiuc.invBoxDiagX);
 159     dr.x    -= shx*pbcAiuc.boxXX;
 160
 161     return dr;
 162 }
 163
 164 #endif //GMX_PBCUTIL_PBC_AIUC_CUDA_CUH