src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel_utils.h

   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 /*! \internal \file
  36  * \brief
  37  * Helper functions and constants for SYCL NBNXM kernels
  38  *
  39  * \ingroup module_nbnxm
  40  */
  41 #ifndef GMX_NBNXM_SYCL_NBNXN_SYCL_KERNEL_UTILS_H
  42 #define GMX_NBNXM_SYCL_NBNXN_SYCL_KERNEL_UTILS_H
  43
  44 #include "gromacs/gpu_utils/sycl_kernel_utils.h"
  45 #include "gromacs/nbnxm/pairlist.h"
  46 #include "gromacs/nbnxm/pairlistparams.h"
  47
  48 namespace Nbnxm
  49 {
  50
  51 #ifndef GMX_NBNXN_PRUNE_KERNEL_J4_CONCURRENCY
  52 #    define GMX_NBNXN_PRUNE_KERNEL_J4_CONCURRENCY 4
  53 #endif
  54 /*! \brief Macro defining default for the prune kernel's j4 processing concurrency.
  55  *
  56  *  The GMX_NBNXN_PRUNE_KERNEL_J4_CONCURRENCY macro allows compile-time override.
  57  */
  58 static constexpr int c_syclPruneKernelJ4Concurrency = GMX_NBNXN_PRUNE_KERNEL_J4_CONCURRENCY;
  59
  60 /* Convenience constants */
  61 /*! \cond */
  62 // cluster size = number of atoms per cluster.
  63 static constexpr int c_clSize = c_nbnxnGpuClusterSize;
  64 // j-cluster size after split (4 in the current implementation).
  65 static constexpr int c_splitClSize = c_clSize / c_nbnxnGpuClusterpairSplit;
  66 // i-cluster interaction mask for a super-cluster with all c_nbnxnGpuNumClusterPerSupercluster=8 bits set.
  67 static constexpr unsigned superClInteractionMask = ((1U << c_nbnxnGpuNumClusterPerSupercluster) - 1U);
  68
  69 // 1/sqrt(pi), same value as \c M_FLOAT_1_SQRTPI in other NB kernels.
  70 static constexpr float c_OneOverSqrtPi = 0.564189583547756F;
  71 // 1/6, same value as in other NB kernels.
  72 static constexpr float c_oneSixth = 0.16666667F;
  73 // 1/12, same value as in other NB kernels.
  74 static constexpr float c_oneTwelfth = 0.08333333F;
  75 /*! \endcond */
  76
  77 /* The following functions are necessary because on some versions of Intel OpenCL RT, subgroups
  78  * do not properly work (segfault or create subgroups of size 1) if used in kernels
  79  * with non-1-dimensional workgroup. */
  80 //! \brief Convert 3D range to 1D
  81 static inline cl::sycl::range<1> flattenRange(cl::sycl::range<3> range3d)
  82 {
  83     return cl::sycl::range<1>(range3d.size());
  84 }
  85
  86 //! \brief Convert 3D nd_range to 1D
  87 static inline cl::sycl::nd_range<1> flattenNDRange(cl::sycl::nd_range<3> nd_range3d)
  88 {
  89     return cl::sycl::nd_range<1>(flattenRange(nd_range3d.get_global_range()),
  90                                  flattenRange(nd_range3d.get_local_range()));
  91 }
  92
  93 //! \brief Convert flattened 1D index to 3D
  94 template<int rangeX, int rangeY>
  95 static inline cl::sycl::id<3> unflattenId(cl::sycl::id<1> id1d)
  96 {
  97     constexpr unsigned rangeXY = rangeX * rangeY;
  98     const unsigned     id      = id1d[0];
  99     const unsigned     z       = id / rangeXY;
 100     const unsigned     xy      = id % rangeXY;
 101     return cl::sycl::id<3>(xy % rangeX, xy / rangeX, z);
 102 }
 103
 104 } // namespace Nbnxm
 105
 106 #endif // GMX_NBNXM_SYCL_NBNXN_SYCL_KERNEL_UTILS_H