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,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.
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/device_stream.h"
45 #include "gromacs/gpu_utils/gputraits.cuh"
46 #include "gromacs/math/vec.h"
47 #include "gromacs/math/vectypes.h"
48 #include "gromacs/utility/exceptions.h"
49 #include "gromacs/utility/fatalerror.h"
50 #include "gromacs/utility/gmxassert.h"
51 #include "gromacs/utility/stringutil.h"
58 /*! \brief Add the API information on the specific error to the error message.
60 * \param[in] deviceError The error to assert cudaSuccess on.
62 * \returns A description of the API error. Returns '(CUDA error #0 (cudaSuccess): no error)' in case deviceError is cudaSuccess.
64 inline std::string getDeviceErrorString(const cudaError_t deviceError)
66 return formatString("CUDA error #%d (%s): %s.",
68 cudaGetErrorName(deviceError),
69 cudaGetErrorString(deviceError));
72 /*! \brief Check if API returned an error and throw an exception with information on it.
74 * \param[in] deviceError The error to assert cudaSuccess on.
75 * \param[in] errorMessage Undecorated error message.
77 * \throws InternalError if deviceError is not a success.
79 inline void checkDeviceError(const cudaError_t deviceError, const std::string& errorMessage)
81 if (deviceError != cudaSuccess)
83 GMX_THROW(gmx::InternalError(errorMessage + " " + getDeviceErrorString(deviceError)));
87 /*! \brief Helper function to ensure no pending error silently
88 * disrupts error handling.
90 * Asserts in a debug build if an unhandled error is present. Issues a
91 * warning at run time otherwise.
93 * \param[in] errorMessage Undecorated error message.
95 inline void ensureNoPendingDeviceError(const std::string& errorMessage)
97 // Ensure there is no pending error that would otherwise affect
98 // the behaviour of future error handling.
99 cudaError_t deviceError = cudaGetLastError();
100 if (deviceError == cudaSuccess)
105 // If we would find an error in a release build, we do not know
106 // what is appropriate to do about it, so assert only for debug
108 const std::string fullErrorMessage =
109 errorMessage + " An unhandled error from a previous CUDA operation was detected. "
110 + gmx::getDeviceErrorString(deviceError);
111 GMX_ASSERT(deviceError == cudaSuccess, fullErrorMessage.c_str());
112 // TODO When we evolve a better logging framework, use that
113 // for release-build error reporting.
114 gmx_warning("%s", fullErrorMessage.c_str());
120 enum class GpuApiCallBehavior;
122 /* TODO error checking needs to be rewritten. We have 2 types of error checks needed
123 based on where they occur in the code:
124 - non performance-critical: these errors are unsafe to be ignored and must be
125 _always_ checked for, e.g. initializations
126 - performance critical: handling errors might hurt performance so care need to be taken
127 when/if we should check for them at all, e.g. in cu_upload_X. However, we should be
128 able to turn the check for these errors on!
130 Probably we'll need two sets of the macros below...
133 #define CHECK_CUDA_ERRORS
135 #ifdef CHECK_CUDA_ERRORS
137 /*! Check for CUDA error on the return status of a CUDA RT API call. */
138 # define CU_RET_ERR(deviceError, msg) \
141 if ((deviceError) != cudaSuccess) \
143 gmx_fatal(FARGS, "%s\n", ((msg) + gmx::getDeviceErrorString(deviceError)).c_str()); \
147 #else /* CHECK_CUDA_ERRORS */
149 # define CU_RET_ERR(status, msg) \
154 #endif /* CHECK_CUDA_ERRORS */
156 // TODO: the 2 functions below are pretty much a constructor/destructor of a simple
157 // GPU table object. There is also almost self-contained fetchFromParamLookupTable()
158 // in cuda_kernel_utils.cuh. They could all live in a separate class/struct file.
160 /*! \brief Add a triplets stored in a float3 to an rvec variable.
162 * \param[out] a Rvec to increment
163 * \param[in] b Float triplet to increment with.
165 static inline void rvec_inc(rvec a, const float3 b)
167 rvec tmp = { b.x, b.y, b.z };
171 /*! \brief Returns true if all tasks in \p s have completed.
173 * \param[in] deviceStream CUDA stream to check.
175 * \returns True if all tasks enqueued in the stream \p deviceStream (at the time of this call) have completed.
177 static inline bool haveStreamTasksCompleted(const DeviceStream& deviceStream)
179 cudaError_t stat = cudaStreamQuery(deviceStream.stream());
181 if (stat == cudaErrorNotReady)
183 // work is still in progress in the stream
187 GMX_ASSERT(stat != cudaErrorInvalidResourceHandle,
188 ("Stream identifier not valid. " + gmx::getDeviceErrorString(stat)).c_str());
190 // cudaSuccess and cudaErrorNotReady are the expected return values
191 CU_RET_ERR(stat, "Unexpected cudaStreamQuery failure. ");
193 GMX_ASSERT(stat == cudaSuccess,
194 ("Values other than cudaSuccess should have been explicitly handled. "
195 + gmx::getDeviceErrorString(stat))
201 /* Kernel launch helpers */
204 * A function for setting up a single CUDA kernel argument.
205 * This is the tail of the compile-time recursive function below.
206 * It has to be seen by the compiler first.
208 * \tparam totalArgsCount Number of the kernel arguments
209 * \tparam KernelPtr Kernel function handle type
210 * \param[in] argIndex Index of the current argument
212 template<size_t totalArgsCount, typename KernelPtr>
213 void prepareGpuKernelArgument(KernelPtr /*kernel*/,
214 std::array<void*, totalArgsCount>* /* kernelArgsPtr */,
215 size_t gmx_used_in_debug argIndex)
217 GMX_ASSERT(argIndex == totalArgsCount, "Tail expansion");
221 * Compile-time recursive function for setting up a single CUDA kernel argument.
222 * This function copies a kernel argument pointer \p argPtr into \p kernelArgsPtr,
223 * and calls itself on the next argument, eventually calling the tail function above.
225 * \tparam CurrentArg Type of the current argument
226 * \tparam RemainingArgs Types of remaining arguments after the current one
227 * \tparam totalArgsCount Number of the kernel arguments
228 * \tparam KernelPtr Kernel function handle type
229 * \param[in] kernel Kernel function handle
230 * \param[in,out] kernelArgsPtr Pointer to the argument array to be filled in
231 * \param[in] argIndex Index of the current argument
232 * \param[in] argPtr Pointer to the current argument
233 * \param[in] otherArgsPtrs Pack of pointers to arguments remaining to process after the current one
235 template<typename CurrentArg, typename... RemainingArgs, size_t totalArgsCount, typename KernelPtr>
236 void prepareGpuKernelArgument(KernelPtr kernel,
237 std::array<void*, totalArgsCount>* kernelArgsPtr,
239 const CurrentArg* argPtr,
240 const RemainingArgs*... otherArgsPtrs)
242 // NOLINTNEXTLINE(google-readability-casting)
243 (*kernelArgsPtr)[argIndex] = (void*)argPtr;
244 prepareGpuKernelArgument(kernel, kernelArgsPtr, argIndex + 1, otherArgsPtrs...);
248 * A wrapper function for setting up all the CUDA kernel arguments.
249 * Calls the recursive functions above.
251 * \tparam KernelPtr Kernel function handle type
252 * \tparam Args Types of all the kernel arguments
253 * \param[in] kernel Kernel function handle
254 * \param[in] argsPtrs Pointers to all the kernel arguments
255 * \returns A prepared parameter pack to be used with launchGpuKernel() as the last argument.
257 template<typename KernelPtr, typename... Args>
258 std::array<void*, sizeof...(Args)> prepareGpuKernelArguments(KernelPtr kernel,
259 const KernelLaunchConfig& /*config */,
260 const Args*... argsPtrs)
262 std::array<void*, sizeof...(Args)> kernelArgs;
263 prepareGpuKernelArgument(kernel, &kernelArgs, 0, argsPtrs...);
267 /*! \brief Launches the CUDA kernel and handles the errors.
269 * \tparam Args Types of all the kernel arguments
270 * \param[in] kernel Kernel function handle
271 * \param[in] config Kernel configuration for launching
272 * \param[in] deviceStream GPU stream to launch kernel in
273 * \param[in] kernelName Human readable kernel description, for error handling only
274 * \param[in] kernelArgs Array of the pointers to the kernel arguments, prepared by
275 * prepareGpuKernelArguments() \throws gmx::InternalError on kernel launch failure
277 template<typename... Args>
278 void launchGpuKernel(void (*kernel)(Args...),
279 const KernelLaunchConfig& config,
280 const DeviceStream& deviceStream,
281 CommandEvent* /*timingEvent */,
282 const char* kernelName,
283 const std::array<void*, sizeof...(Args)>& kernelArgs)
285 dim3 blockSize(config.blockSize[0], config.blockSize[1], config.blockSize[2]);
286 dim3 gridSize(config.gridSize[0], config.gridSize[1], config.gridSize[2]);
287 // NOLINTNEXTLINE(google-readability-casting)
288 cudaLaunchKernel((void*)kernel,
291 const_cast<void**>(kernelArgs.data()),
292 config.sharedMemorySize,
293 deviceStream.stream());
295 gmx::ensureNoPendingDeviceError("GPU kernel (" + std::string(kernelName)
296 + ") failed to launch.");