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