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