src/gromacs/gpu_utils/cuda_arch_utils.cuh

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2012,2014,2015,2016,2017 by the GROMACS development team.
   5  * Copyright (c) 2018,2019,2020,2021, by the GROMACS development team, led by
   6  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   7  * and including many others, as listed in the AUTHORS file in the
   8  * top-level source directory and at http://www.gromacs.org.
   9  *
  10  * GROMACS is free software; you can redistribute it and/or
  11  * modify it under the terms of the GNU Lesser General Public License
  12  * as published by the Free Software Foundation; either version 2.1
  13  * of the License, or (at your option) any later version.
  14  *
  15  * GROMACS is distributed in the hope that it will be useful,
  16  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  17  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  18  * Lesser General Public License for more details.
  19  *
  20  * You should have received a copy of the GNU Lesser General Public
  21  * License along with GROMACS; if not, see
  22  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  23  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  24  *
  25  * If you want to redistribute modifications to GROMACS, please
  26  * consider that scientific software is very special. Version
  27  * control is crucial - bugs must be traceable. We will be happy to
  28  * consider code for inclusion in the official distribution, but
  29  * derived work must not be called official GROMACS. Details are found
  30  * in the README & COPYING files - if they are missing, get the
  31  * official version at http://www.gromacs.org.
  32  *
  33  * To help us fund GROMACS development, we humbly ask that you cite
  34  * the research papers on the package. Check out http://www.gromacs.org.
  35  */
  36 #ifndef CUDA_ARCH_UTILS_CUH_
  37 #define CUDA_ARCH_UTILS_CUH_
  38
  39 #include "gromacs/utility/basedefinitions.h"
  40
  41 /*! \file
  42  *  \brief CUDA arch dependent definitions.
  43  *
  44  *  \author Szilard Pall <pall.szilard@gmail.com>
  45  */
  46
  47 /* GMX_PTX_ARCH is set to the virtual arch (PTX) version targeted by
  48  * the current compiler pass or zero for the host pass and it is
  49  * intended to be used instead of __CUDA_ARCH__.
  50  */
  51 #ifndef __CUDA_ARCH__
  52 #    define GMX_PTX_ARCH 0
  53 #else
  54 #    define GMX_PTX_ARCH __CUDA_ARCH__
  55 #endif
  56
  57 /* Until CC 5.2 and likely for the near future all NVIDIA architectures
  58    have a warp size of 32, but this could change later. If it does, the
  59    following constants should depend on the value of GMX_PTX_ARCH.
  60  */
  61 static const int warp_size      = 32;
  62 static const int warp_size_log2 = 5;
  63 /*! \brief Bitmask corresponding to all threads active in a warp.
  64  *  NOTE that here too we assume 32-wide warps.
  65  */
  66 static const unsigned int c_fullWarpMask = 0xffffffff;
  67
  68 /*! \brief Allow disabling CUDA textures using the GMX_DISABLE_CUDA_TEXTURES macro.
  69  *
  70  *  Only texture objects supported.
  71  *  Disable texture support missing in clang (all versions up to <=5.0-dev as of writing).
  72  *  Disable texture support on CC 7.0 and 8.0 for performance reasons (Issue #3845).
  73  *
  74  *  This option will not influence functionality. All features using textures ought
  75  *  to have fallback for texture-less reads (direct/LDG loads), all new code needs
  76  *  to provide fallback code.
  77  */
  78 #if defined(GMX_DISABLE_CUDA_TEXTURES) || (defined(__clang__) && defined(__CUDA__)) \
  79         || (GMX_PTX_ARCH == 700) || (GMX_PTX_ARCH == 800)
  80 #    define DISABLE_CUDA_TEXTURES 1
  81 #else
  82 #    define DISABLE_CUDA_TEXTURES 0
  83 #endif
  84
  85 /*! \brief True if the use of texture fetch in the CUDA kernels is disabled. */
  86 static const bool c_disableCudaTextures = DISABLE_CUDA_TEXTURES;
  87
  88
  89 /* CUDA architecture technical characteristics. Needs macros because it is used
  90  * in the __launch_bounds__ function qualifiers and might need it in preprocessor
  91  * conditionals.
  92  *
  93  */
  94 #if GMX_PTX_ARCH > 0
  95 #    if GMX_PTX_ARCH <= 370 // CC 3.x
  96 #        define GMX_CUDA_MAX_BLOCKS_PER_MP 16
  97 #        define GMX_CUDA_MAX_THREADS_PER_MP 2048
  98 #    elif GMX_PTX_ARCH == 750 // CC 7.5, lower limits compared to 7.0
  99 #        define GMX_CUDA_MAX_BLOCKS_PER_MP 16
 100 #        define GMX_CUDA_MAX_THREADS_PER_MP 1024
 101 #    elif GMX_PTX_ARCH == 860 // CC 8.6, lower limits compared to 8.0
 102 #        define GMX_CUDA_MAX_BLOCKS_PER_MP 16
 103 #        define GMX_CUDA_MAX_THREADS_PER_MP 1536
 104 #    else // CC 5.x, 6.x, 7.0, 8.0
 105 /* Note that this final branch covers all future architectures (current gen
 106  * is 8.x as of writing), hence assuming that these *currently defined* upper
 107  * limits will not be lowered.
 108  */
 109 #        define GMX_CUDA_MAX_BLOCKS_PER_MP 32
 110 #        define GMX_CUDA_MAX_THREADS_PER_MP 2048
 111 #    endif
 112 #else
 113 #    define GMX_CUDA_MAX_BLOCKS_PER_MP 0
 114 #    define GMX_CUDA_MAX_THREADS_PER_MP 0
 115 #endif
 116
 117 // Macro defined for clang CUDA device compilation in the presence of debug symbols
 118 // used to work around codegen bug that breaks some kernels when assertions are on
 119 // at -O1 and higher (tested with clang 6-8).
 120 #if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__) && !defined(NDEBUG)
 121 #    define CLANG_DISABLE_OPTIMIZATION_ATTRIBUTE __attribute__((optnone))
 122 #else
 123 #    define CLANG_DISABLE_OPTIMIZATION_ATTRIBUTE
 124 #endif
 125
 126
 127 #endif /* CUDA_ARCH_UTILS_CUH_ */