2 * This file is part of the GROMACS molecular simulation package.
4 * Copyright (c) 2012,2014,2015,2016,2017,2018, 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.
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.
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.
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.
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.
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.
35 #ifndef GMX_GPU_UTILS_CUDAUTILS_CUH
36 #define GMX_GPU_UTILS_CUDAUTILS_CUH
43 #endif /* HAVE_NVML */
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"
58 /*! \brief Helper function to ensure no pending error silently
59 * disrupts error handling.
61 * Asserts in a debug build if an unhandled error is present. Issues a
62 * warning at run time otherwise.
64 * \todo This is similar to CU_CHECK_PREV_ERR, which should be
67 static inline void ensureNoPendingCudaError(const char *errorMessage)
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)
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
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());
91 enum class GpuApiCallBehavior;
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!
101 Probably we'll need two sets of the macros below...
104 #define CHECK_CUDA_ERRORS
106 #ifdef CHECK_CUDA_ERRORS
108 /*! Check for CUDA error on the return status of a CUDA RT API call. */
109 #define CU_RET_ERR(status, msg) \
111 if (status != cudaSuccess) \
113 gmx_fatal(FARGS, "%s: %s\n", msg, cudaGetErrorString(status)); \
117 /*! Check for any previously occurred uncaught CUDA error. */
118 #define CU_CHECK_PREV_ERR() \
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)); \
126 /*! Check for any previously occurred uncaught CUDA error
127 -- aimed at use after kernel calls. */
128 #define CU_LAUNCH_ERR(msg) \
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)); \
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) \
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)); \
146 #else /* CHECK_CUDA_ERRORS */
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)
154 #endif /* CHECK_CUDA_ERRORS */
156 /*! \brief CUDA device information.
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.
162 * \todo extract an object to manage NVML details
164 struct gmx_device_info_t
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 */
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 */
181 /*! Launches synchronous or asynchronous device to host memory copy.
183 * The copy is launched in stream s or if not specified, in stream 0.
185 int cu_copy_D2H(void *h_dest, void *d_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t s /*= 0*/);
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*/);
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*/);
193 /*! Launches synchronous or asynchronous host to device memory copy.
195 * The copy is launched in stream s or if not specified, in stream 0.
197 int cu_copy_H2D(void *d_dest, void *h_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t /*s = 0*/);
199 /*! Launches synchronous host to device memory copy. */
200 int cu_copy_H2D_sync(void * /*d_dest*/, void * /*h_src*/, size_t /*bytes*/);
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*/);
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);
208 /*! Reallocates the device memory and copies data from the host. */
209 void cu_realloc_buffered(void **d_dest, void *h_src,
211 int *curr_size, int *curr_alloc_size,
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.
220 /*! \brief Initialize parameter lookup table.
222 * Initializes device memory, copies data from host and binds
223 * a texture to allocated device memory to be used for parameter lookup.
225 * \tparam[in] T Raw data type
226 * \param[out] d_ptr device pointer to the memory to be allocated
227 * \param[out] texObj texture object to be initialized
228 * \param[in] h_ptr pointer to the host memory to be uploaded to the device
229 * \param[in] numElem number of elements in the h_ptr
230 * \param[in] devInfo pointer to the info struct of the device in use
232 template <typename T>
233 void initParamLookupTable(T * &d_ptr,
234 cudaTextureObject_t &texObj,
237 const gmx_device_info_t *devInfo);
239 /*! \brief Destroy parameter lookup table.
241 * Unbinds texture object, deallocates device memory.
243 * \tparam[in] T Raw data type
244 * \param[in] d_ptr Device pointer to the memory to be deallocated
245 * \param[in] texObj Texture object to be deinitialized
246 * \param[in] devInfo Pointer to the info struct of the device in use
248 template <typename T>
249 void destroyParamLookupTable(T *d_ptr,
250 cudaTextureObject_t texObj,
251 const gmx_device_info_t *devInfo);
253 /*! \brief Add a triplets stored in a float3 to an rvec variable.
255 * \param[out] a Rvec to increment
256 * \param[in] b Float triplet to increment with.
258 static inline void rvec_inc(rvec a, const float3 b)
260 rvec tmp = {b.x, b.y, b.z};
264 /*! \brief Wait for all taks in stream \p s to complete.
266 * \param[in] s stream to synchronize with
268 static inline void gpuStreamSynchronize(cudaStream_t s)
270 cudaError_t stat = cudaStreamSynchronize(s);
271 CU_RET_ERR(stat, "cudaStreamSynchronize failed");
274 /*! \brief Returns true if all tasks in \p s have completed.
276 * \param[in] s stream to check
278 * \returns True if all tasks enqueued in the stream \p s (at the time of this call) have completed.
280 static inline bool haveStreamTasksCompleted(cudaStream_t s)
282 cudaError_t stat = cudaStreamQuery(s);
284 if (stat == cudaErrorNotReady)
286 // work is still in progress in the stream
290 GMX_ASSERT(stat != cudaErrorInvalidResourceHandle, "Stream idnetifier not valid");
292 // cudaSuccess and cudaErrorNotReady are the expected return values
293 CU_RET_ERR(stat, "Unexpected cudaStreamQuery failure");
295 GMX_ASSERT(stat == cudaSuccess, "Values other than cudaSuccess should have been explicitly handled");
300 /*! \brief Free a device-side buffer.
301 * TODO: fully replace cu_free_buffered with this.
303 * \param[in] buffer Pointer to the buffer to free.
305 template <typename DeviceBuffer>
306 void freeDeviceBuffer(DeviceBuffer *buffer)
308 GMX_ASSERT(buffer, "needs a buffer pointer");
311 GMX_RELEASE_ASSERT(cudaFree(*buffer) == cudaSuccess, "cudaFree failed");