src/gromacs/gpu_utils/cudautils.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, 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 #ifndef GMX_GPU_UTILS_CUDAUTILS_CUH
  36 #define GMX_GPU_UTILS_CUDAUTILS_CUH
  37
  38 #include "config.h"
  39
  40 #include <stdio.h>
  41 #if HAVE_NVML
  42 #include <nvml.h>
  43 #endif /* HAVE_NVML */
  44
  45 #include <string>
  46
  47 #include "gromacs/math/vec.h"
  48 #include "gromacs/math/vectypes.h"
  49 #include "gromacs/utility/fatalerror.h"
  50 #include "gromacs/utility/gmxassert.h"
  51 #include "gromacs/utility/stringutil.h"
  52
  53 namespace gmx
  54 {
  55 namespace
  56 {
  57
  58 /*! \brief Helper function to ensure no pending error silently
  59  * disrupts error handling.
  60  *
  61  * Asserts in a debug build if an unhandled error is present. Issues a
  62  * warning at run time otherwise.
  63  *
  64  * \todo This is similar to CU_CHECK_PREV_ERR, which should be
  65  * consolidated.
  66  */
  67 static inline void ensureNoPendingCudaError(const char *errorMessage)
  68 {
  69     // Ensure there is no pending error that would otherwise affect
  70     // the behaviour of future error handling.
  71     cudaError_t stat = cudaGetLastError();
  72     if (stat == cudaSuccess)
  73     {
  74         return;
  75     }
  76
  77     // If we would find an error in a release build, we do not know
  78     // what is appropriate to do about it, so assert only for debug
  79     // builds.
  80     auto fullMessage = formatString("%s An unhandled error from a previous CUDA operation was detected. %s: %s",
  81                                     errorMessage, cudaGetErrorName(stat), cudaGetErrorString(stat));
  82     GMX_ASSERT(stat == cudaSuccess, fullMessage.c_str());
  83     // TODO When we evolve a better logging framework, use that
  84     // for release-build error reporting.
  85     gmx_warning(fullMessage.c_str());
  86 }
  87
  88 }   // namespace
  89 }   // namespace
  90
  91 enum class GpuApiCallBehavior;
  92
  93 /* TODO error checking needs to be rewritten. We have 2 types of error checks needed
  94    based on where they occur in the code:
  95    - non performance-critical: these errors are unsafe to be ignored and must be
  96      _always_ checked for, e.g. initializations
  97    - performance critical: handling errors might hurt performance so care need to be taken
  98      when/if we should check for them at all, e.g. in cu_upload_X. However, we should be
  99      able to turn the check for these errors on!
 100
 101    Probably we'll need two sets of the macros below...
 102
 103  */
 104 #define CHECK_CUDA_ERRORS
 105
 106 #ifdef CHECK_CUDA_ERRORS
 107
 108 /*! Check for CUDA error on the return status of a CUDA RT API call. */
 109 #define CU_RET_ERR(status, msg) \
 110     do { \
 111         if (status != cudaSuccess) \
 112         { \
 113             gmx_fatal(FARGS, "%s: %s\n", msg, cudaGetErrorString(status)); \
 114         } \
 115     } while (0)
 116
 117 /*! Check for any previously occurred uncaught CUDA error. */
 118 #define CU_CHECK_PREV_ERR() \
 119     do { \
 120         cudaError_t _CU_CHECK_PREV_ERR_status = cudaGetLastError(); \
 121         if (_CU_CHECK_PREV_ERR_status != cudaSuccess) { \
 122             gmx_warning("Just caught a previously occurred CUDA error (%s), will try to continue.", cudaGetErrorString(_CU_CHECK_PREV_ERR_status)); \
 123         } \
 124     } while (0)
 125
 126 /*! Check for any previously occurred uncaught CUDA error
 127    -- aimed at use after kernel calls. */
 128 #define CU_LAUNCH_ERR(msg) \
 129     do { \
 130         cudaError_t _CU_LAUNCH_ERR_status = cudaGetLastError(); \
 131         if (_CU_LAUNCH_ERR_status != cudaSuccess) { \
 132             gmx_fatal(FARGS, "Error while launching kernel %s: %s\n", msg, cudaGetErrorString(_CU_LAUNCH_ERR_status)); \
 133         } \
 134     } while (0)
 135
 136 /*! Synchronize with GPU and check for any previously occurred uncaught CUDA error
 137    -- aimed at use after kernel calls. */
 138 #define CU_LAUNCH_ERR_SYNC(msg) \
 139     do { \
 140         cudaError_t _CU_SYNC_LAUNCH_ERR_status = cudaThreadSynchronize(); \
 141         if (_CU_SYNC_LAUNCH_ERR_status != cudaSuccess) { \
 142             gmx_fatal(FARGS, "Error while launching kernel %s: %s\n", msg, cudaGetErrorString(_CU_SYNC_LAUNCH_ERR_status)); \
 143         } \
 144     } while (0)
 145
 146 #else /* CHECK_CUDA_ERRORS */
 147
 148 #define CU_RET_ERR(status, msg) do { } while (0)
 149 #define CU_CHECK_PREV_ERR()     do { } while (0)
 150 #define CU_LAUNCH_ERR(msg)      do { } while (0)
 151 #define CU_LAUNCH_ERR_SYNC(msg) do { } while (0)
 152 #define HANDLE_NVML_RET_ERR(status, msg) do { } while (0)
 153
 154 #endif /* CHECK_CUDA_ERRORS */
 155
 156 /*! \brief CUDA device information.
 157  *
 158  * The CUDA device information is queried and set at detection and contains
 159  * both information about the device/hardware returned by the runtime as well
 160  * as additional data like support status.
 161  *
 162  * \todo extract an object to manage NVML details
 163  */
 164 struct gmx_device_info_t
 165 {
 166     int                 id;                      /* id of the CUDA device */
 167     cudaDeviceProp      prop;                    /* CUDA device properties */
 168     int                 stat;                    /* result of the device check */
 169     unsigned int        nvml_orig_app_sm_clock;  /* The original SM clock before we changed it */
 170     unsigned int        nvml_orig_app_mem_clock; /* The original memory clock before we changed it */
 171     gmx_bool            nvml_app_clocks_changed; /* If application clocks have been changed */
 172     unsigned int        nvml_set_app_sm_clock;   /* The SM clock we set */
 173     unsigned int        nvml_set_app_mem_clock;  /* The memory clock we set */
 174 #if HAVE_NVML
 175     nvmlDevice_t        nvml_device_id;          /* NVML device id */
 176     // TODO This can become a bool with a more useful name
 177     nvmlEnableState_t   nvml_is_restricted;      /* Status of application clocks permission */
 178 #endif                                           /* HAVE_NVML */
 179 };
 180
 181 /*! Launches synchronous or asynchronous device to host memory copy.
 182  *
 183  *  The copy is launched in stream s or if not specified, in stream 0.
 184  */
 185 int cu_copy_D2H(void *h_dest, void *d_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t s /*= 0*/);
 186
 187 /*! Launches synchronous host to device memory copy in stream 0. */
 188 int cu_copy_D2H_sync(void * /*h_dest*/, void * /*d_src*/, size_t /*bytes*/);
 189
 190 /*! Launches asynchronous host to device memory copy in stream s. */
 191 int cu_copy_D2H_async(void * /*h_dest*/, void * /*d_src*/, size_t /*bytes*/, cudaStream_t /*s = 0*/);
 192
 193 /*! Launches synchronous or asynchronous host to device memory copy.
 194  *
 195  *  The copy is launched in stream s or if not specified, in stream 0.
 196  */
 197 int cu_copy_H2D(void *d_dest, void *h_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t /*s = 0*/);
 198
 199 /*! Launches synchronous host to device memory copy. */
 200 int cu_copy_H2D_sync(void * /*d_dest*/, void * /*h_src*/, size_t /*bytes*/);
 201
 202 /*! Launches asynchronous host to device memory copy in stream s. */
 203 int cu_copy_H2D_async(void * /*d_dest*/, void * /*h_src*/, size_t /*bytes*/, cudaStream_t /*s = 0*/);
 204
 205 /*! Frees device memory and resets the size and allocation size to -1. */
 206 void cu_free_buffered(void *d_ptr, int *n = NULL, int *nalloc = NULL);
 207
 208 /*! Reallocates the device memory and copies data from the host. */
 209 void cu_realloc_buffered(void **d_dest, void *h_src,
 210                          size_t type_size,
 211                          int *curr_size, int *curr_alloc_size,
 212                          int req_size,
 213                          cudaStream_t s,
 214                          bool bAsync);
 215
 216 // TODO: the 2 functions below are pretty much a constructor/destructor of a simple
 217 // GPU table object. There is also almost self-contained fetchFromParamLookupTable()
 218 // in cuda_kernel_utils.cuh. They could all live in a separate class/struct file,
 219 // granted storing static texture references in there does not pose problems.
 220
 221 /*! \brief Initialize parameter lookup table.
 222  *
 223  * Initializes device memory, copies data from host and binds
 224  * a texture to allocated device memory to be used for parameter lookup.
 225  *
 226  * \tparam[in] T         Raw data type
 227  * \param[out] d_ptr     device pointer to the memory to be allocated
 228  * \param[out] texObj    texture object to be initialized
 229  * \param[out] texRef    texture reference to be initialized
 230  * \param[in]  h_ptr     pointer to the host memory to be uploaded to the device
 231  * \param[in]  numElem   number of elements in the h_ptr
 232  * \param[in]  devInfo   pointer to the info struct of the device in use
 233  */
 234 template <typename T>
 235 void initParamLookupTable(T                        * &d_ptr,
 236                           cudaTextureObject_t       &texObj,
 237                           const struct texture<T, 1, cudaReadModeElementType> *texRef,
 238                           const T                   *h_ptr,
 239                           int                        numElem,
 240                           const gmx_device_info_t   *devInfo);
 241
 242 /*! \brief Destroy parameter lookup table.
 243  *
 244  * Unbinds texture reference/object, deallocates device memory.
 245  *
 246  * \tparam[in] T         Raw data type
 247  * \param[in]  d_ptr     Device pointer to the memory to be deallocated
 248  * \param[in]  texObj    Texture object to be deinitialized
 249  * \param[in]  texRef    Texture reference to be deinitialized
 250  * \param[in]  devInfo   Pointer to the info struct of the device in use
 251  */
 252 template <typename T>
 253 void destroyParamLookupTable(T                         *d_ptr,
 254                              cudaTextureObject_t        texObj,
 255                              const struct texture<T, 1, cudaReadModeElementType> *texRef,
 256                              const gmx_device_info_t   *devInfo);
 257
 258 /*! \brief Add a triplets stored in a float3 to an rvec variable.
 259  *
 260  * \param[out]  a Rvec to increment
 261  * \param[in]   b Float triplet to increment with.
 262  */
 263 static inline void rvec_inc(rvec a, const float3 b)
 264 {
 265     rvec tmp = {b.x, b.y, b.z};
 266     rvec_inc(a, tmp);
 267 }
 268
 269 /*! \brief Calls cudaStreamSynchronize() in the stream \p s.
 270  *
 271  * \param[in] s stream to synchronize with
 272  */
 273 static inline void gpuStreamSynchronize(cudaStream_t s)
 274 {
 275     cudaError_t stat = cudaStreamSynchronize(s);
 276     CU_RET_ERR(stat, "cudaStreamSynchronize failed");
 277 }
 278
 279 #endif