2 * This file is part of the GROMACS molecular simulation package.
4 * Copyright (c) 2012,2014,2015,2016,2017 by the GROMACS development team.
5 * Copyright (c) 2018,2019,2020, 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.
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.
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.
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.
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.
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.
36 #ifndef GMX_GPU_UTILS_CUDAUTILS_CUH
37 #define GMX_GPU_UTILS_CUDAUTILS_CUH
44 #include "gromacs/gpu_utils/gputraits.cuh"
45 #include "gromacs/math/vec.h"
46 #include "gromacs/math/vectypes.h"
47 #include "gromacs/utility/exceptions.h"
48 #include "gromacs/utility/fatalerror.h"
49 #include "gromacs/utility/gmxassert.h"
50 #include "gromacs/utility/stringutil.h"
57 /*! \brief Helper function to ensure no pending error silently
58 * disrupts error handling.
60 * Asserts in a debug build if an unhandled error is present. Issues a
61 * warning at run time otherwise.
63 * \todo This is similar to CU_CHECK_PREV_ERR, which should be
66 static inline void ensureNoPendingCudaError(const char* errorMessage)
68 // Ensure there is no pending error that would otherwise affect
69 // the behaviour of future error handling.
70 cudaError_t stat = cudaGetLastError();
71 if (stat == cudaSuccess)
76 // If we would find an error in a release build, we do not know
77 // what is appropriate to do about it, so assert only for debug
79 auto fullMessage = formatString(
80 "%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("%s", 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) \
112 if (status != cudaSuccess) \
114 gmx_fatal(FARGS, "%s: %s\n", msg, cudaGetErrorString(status)); \
118 /*! Check for any previously occurred uncaught CUDA error. */
119 # define CU_CHECK_PREV_ERR() \
122 cudaError_t _CU_CHECK_PREV_ERR_status = cudaGetLastError(); \
123 if (_CU_CHECK_PREV_ERR_status != cudaSuccess) \
126 "Just caught a previously occurred CUDA error (%s), will try to " \
128 cudaGetErrorString(_CU_CHECK_PREV_ERR_status)); \
132 #else /* CHECK_CUDA_ERRORS */
134 # define CU_RET_ERR(status, msg) \
138 # define CU_CHECK_PREV_ERR() \
143 #endif /* CHECK_CUDA_ERRORS */
145 /*! Launches synchronous or asynchronous device to host memory copy.
147 * The copy is launched in stream s or if not specified, in stream 0.
149 int cu_copy_D2H(void* h_dest, void* d_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t /*s = nullptr*/);
151 /*! Launches synchronous host to device memory copy in stream 0. */
152 int cu_copy_D2H_sync(void* /*h_dest*/, void* /*d_src*/, size_t /*bytes*/);
154 /*! Launches asynchronous host to device memory copy in stream s. */
155 int cu_copy_D2H_async(void* /*h_dest*/, void* /*d_src*/, size_t /*bytes*/, cudaStream_t /*s = nullptr*/);
157 /*! Launches synchronous or asynchronous host to device memory copy.
159 * The copy is launched in stream s or if not specified, in stream 0.
161 int cu_copy_H2D(void* d_dest, const void* h_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t /*s = nullptr*/);
163 /*! Launches synchronous host to device memory copy. */
164 int cu_copy_H2D_sync(void* /*d_dest*/, const void* /*h_src*/, size_t /*bytes*/);
166 /*! Launches asynchronous host to device memory copy in stream s. */
167 int cu_copy_H2D_async(void* /*d_dest*/, const void* /*h_src*/, size_t /*bytes*/, cudaStream_t /*s = nullptr*/);
169 // TODO: the 2 functions below are pretty much a constructor/destructor of a simple
170 // GPU table object. There is also almost self-contained fetchFromParamLookupTable()
171 // in cuda_kernel_utils.cuh. They could all live in a separate class/struct file.
173 /*! \brief Initialize parameter lookup table.
175 * Initializes device memory, copies data from host and binds
176 * a texture to allocated device memory to be used for parameter lookup.
178 * \tparam[in] T Raw data type
179 * \param[out] d_ptr device pointer to the memory to be allocated
180 * \param[out] texObj texture object to be initialized
181 * \param[in] h_ptr pointer to the host memory to be uploaded to the device
182 * \param[in] numElem number of elements in the h_ptr
185 void initParamLookupTable(T*& d_ptr, cudaTextureObject_t& texObj, const T* h_ptr, int numElem);
187 // Add extern declarations so each translation unit understands that
188 // there will be a definition provided.
189 extern template void initParamLookupTable<int>(int*&, cudaTextureObject_t&, const int*, int);
190 extern template void initParamLookupTable<float>(float*&, cudaTextureObject_t&, const float*, int);
192 /*! \brief Destroy parameter lookup table.
194 * Unbinds texture object, deallocates device memory.
196 * \tparam[in] T Raw data type
197 * \param[in] d_ptr Device pointer to the memory to be deallocated
198 * \param[in] texObj Texture object to be deinitialized
201 void destroyParamLookupTable(T* d_ptr, cudaTextureObject_t texObj);
203 // Add extern declarations so each translation unit understands that
204 // there will be a definition provided.
205 extern template void destroyParamLookupTable<int>(int*, cudaTextureObject_t);
206 extern template void destroyParamLookupTable<float>(float*, cudaTextureObject_t);
208 /*! \brief Add a triplets stored in a float3 to an rvec variable.
210 * \param[out] a Rvec to increment
211 * \param[in] b Float triplet to increment with.
213 static inline void rvec_inc(rvec a, const float3 b)
215 rvec tmp = { b.x, b.y, b.z };
219 /*! \brief Wait for all taks in stream \p s to complete.
221 * \param[in] s stream to synchronize with
223 static inline void gpuStreamSynchronize(cudaStream_t s)
225 cudaError_t stat = cudaStreamSynchronize(s);
226 CU_RET_ERR(stat, "cudaStreamSynchronize failed");
229 /*! \brief Returns true if all tasks in \p s have completed.
231 * \param[in] s stream to check
233 * \returns True if all tasks enqueued in the stream \p s (at the time of this call) have completed.
235 static inline bool haveStreamTasksCompleted(cudaStream_t s)
237 cudaError_t stat = cudaStreamQuery(s);
239 if (stat == cudaErrorNotReady)
241 // work is still in progress in the stream
245 GMX_ASSERT(stat != cudaErrorInvalidResourceHandle, "Stream idnetifier not valid");
247 // cudaSuccess and cudaErrorNotReady are the expected return values
248 CU_RET_ERR(stat, "Unexpected cudaStreamQuery failure");
250 GMX_ASSERT(stat == cudaSuccess,
251 "Values other than cudaSuccess should have been explicitly handled");
256 /* Kernel launch helpers */
259 * A function for setting up a single CUDA kernel argument.
260 * This is the tail of the compile-time recursive function below.
261 * It has to be seen by the compiler first.
263 * \tparam totalArgsCount Number of the kernel arguments
264 * \tparam KernelPtr Kernel function handle type
265 * \param[in] argIndex Index of the current argument
267 template<size_t totalArgsCount, typename KernelPtr>
268 void prepareGpuKernelArgument(KernelPtr /*kernel*/,
269 std::array<void*, totalArgsCount>* /* kernelArgsPtr */,
270 size_t gmx_used_in_debug argIndex)
272 GMX_ASSERT(argIndex == totalArgsCount, "Tail expansion");
276 * Compile-time recursive function for setting up a single CUDA kernel argument.
277 * This function copies a kernel argument pointer \p argPtr into \p kernelArgsPtr,
278 * and calls itself on the next argument, eventually calling the tail function above.
280 * \tparam CurrentArg Type of the current argument
281 * \tparam RemainingArgs Types of remaining arguments after the current one
282 * \tparam totalArgsCount Number of the kernel arguments
283 * \tparam KernelPtr Kernel function handle type
284 * \param[in] kernel Kernel function handle
285 * \param[in,out] kernelArgsPtr Pointer to the argument array to be filled in
286 * \param[in] argIndex Index of the current argument
287 * \param[in] argPtr Pointer to the current argument
288 * \param[in] otherArgsPtrs Pack of pointers to arguments remaining to process after the current one
290 template<typename CurrentArg, typename... RemainingArgs, size_t totalArgsCount, typename KernelPtr>
291 void prepareGpuKernelArgument(KernelPtr kernel,
292 std::array<void*, totalArgsCount>* kernelArgsPtr,
294 const CurrentArg* argPtr,
295 const RemainingArgs*... otherArgsPtrs)
297 (*kernelArgsPtr)[argIndex] = (void*)argPtr;
298 prepareGpuKernelArgument(kernel, kernelArgsPtr, argIndex + 1, otherArgsPtrs...);
302 * A wrapper function for setting up all the CUDA kernel arguments.
303 * Calls the recursive functions above.
305 * \tparam KernelPtr Kernel function handle type
306 * \tparam Args Types of all the kernel arguments
307 * \param[in] kernel Kernel function handle
308 * \param[in] argsPtrs Pointers to all the kernel arguments
309 * \returns A prepared parameter pack to be used with launchGpuKernel() as the last argument.
311 template<typename KernelPtr, typename... Args>
312 std::array<void*, sizeof...(Args)> prepareGpuKernelArguments(KernelPtr kernel,
313 const KernelLaunchConfig& /*config */,
314 const Args*... argsPtrs)
316 std::array<void*, sizeof...(Args)> kernelArgs;
317 prepareGpuKernelArgument(kernel, &kernelArgs, 0, argsPtrs...);
321 /*! \brief Launches the CUDA kernel and handles the errors.
323 * \tparam Args Types of all the kernel arguments
324 * \param[in] kernel Kernel function handle
325 * \param[in] config Kernel configuration for launching
326 * \param[in] kernelName Human readable kernel description, for error handling only
327 * \param[in] kernelArgs Array of the pointers to the kernel arguments, prepared by
328 * prepareGpuKernelArguments() \throws gmx::InternalError on kernel launch failure
330 template<typename... Args>
331 void launchGpuKernel(void (*kernel)(Args...),
332 const KernelLaunchConfig& config,
333 CommandEvent* /*timingEvent */,
334 const char* kernelName,
335 const std::array<void*, sizeof...(Args)>& kernelArgs)
337 dim3 blockSize(config.blockSize[0], config.blockSize[1], config.blockSize[2]);
338 dim3 gridSize(config.gridSize[0], config.gridSize[1], config.gridSize[2]);
339 cudaLaunchKernel((void*)kernel, gridSize, blockSize, const_cast<void**>(kernelArgs.data()),
340 config.sharedMemorySize, config.stream);
342 cudaError_t status = cudaGetLastError();
343 if (cudaSuccess != status)
345 const std::string errorMessage =
346 "GPU kernel (" + std::string(kernelName)
347 + ") failed to launch: " + std::string(cudaGetErrorString(status));
348 GMX_THROW(gmx::InternalError(errorMessage));