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.
   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.
   9  *
  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.
  14  *
  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.
  19  *
  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.
  24  *
  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.
  32  *
  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.
  35  */
  36 #ifndef GMX_GPU_UTILS_CUDAUTILS_CUH
  37 #define GMX_GPU_UTILS_CUDAUTILS_CUH
  38
  39 #include <stdio.h>
  40
  41 #include <array>
  42 #include <string>
  43 #include <type_traits>
  44
  45 #include "gromacs/gpu_utils/device_stream.h"
  46 #include "gromacs/gpu_utils/gputraits.cuh"
  47 #include "gromacs/math/vec.h"
  48 #include "gromacs/math/vectypes.h"
  49 #include "gromacs/utility/exceptions.h"
  50 #include "gromacs/utility/fatalerror.h"
  51 #include "gromacs/utility/gmxassert.h"
  52 #include "gromacs/utility/stringutil.h"
  53
  54 namespace gmx
  55 {
  56 namespace
  57 {
  58
  59 /*! \brief Add the API information on the specific error to the error message.
  60  *
  61  * \param[in]  deviceError  The error to assert cudaSuccess on.
  62  *
  63  * \returns A description of the API error. Returns '(CUDA error #0 (cudaSuccess): no error)' in case deviceError is cudaSuccess.
  64  */
  65 inline std::string getDeviceErrorString(const cudaError_t deviceError)
  66 {
  67     return formatString("CUDA error #%d (%s): %s.",
  68                         deviceError,
  69                         cudaGetErrorName(deviceError),
  70                         cudaGetErrorString(deviceError));
  71 }
  72
  73 /*! \brief Check if API returned an error and throw an exception with information on it.
  74  *
  75  * \param[in]  deviceError  The error to assert cudaSuccess on.
  76  * \param[in]  errorMessage  Undecorated error message.
  77  *
  78  *  \throws InternalError if deviceError is not a success.
  79  */
  80 inline void checkDeviceError(const cudaError_t deviceError, const std::string& errorMessage)
  81 {
  82     if (deviceError != cudaSuccess)
  83     {
  84         GMX_THROW(gmx::InternalError(errorMessage + " " + getDeviceErrorString(deviceError)));
  85     }
  86 }
  87
  88 /*! \brief Helper function to ensure no pending error silently
  89  * disrupts error handling.
  90  *
  91  * Asserts in a debug build if an unhandled error is present. Issues a
  92  * warning at run time otherwise.
  93  *
  94  * \param[in]  errorMessage  Undecorated error message.
  95  */
  96 inline void ensureNoPendingDeviceError(const std::string& errorMessage)
  97 {
  98     // Ensure there is no pending error that would otherwise affect
  99     // the behaviour of future error handling.
 100     cudaError_t deviceError = cudaGetLastError();
 101     if (deviceError == cudaSuccess)
 102     {
 103         return;
 104     }
 105
 106     // If we would find an error in a release build, we do not know
 107     // what is appropriate to do about it, so assert only for debug
 108     // builds.
 109     const std::string fullErrorMessage =
 110             errorMessage + " An unhandled error from a previous CUDA operation was detected. "
 111             + gmx::getDeviceErrorString(deviceError);
 112     GMX_ASSERT(deviceError == cudaSuccess, fullErrorMessage.c_str());
 113     // TODO When we evolve a better logging framework, use that
 114     // for release-build error reporting.
 115     gmx_warning("%s", fullErrorMessage.c_str());
 116 }
 117
 118 } // namespace
 119 } // namespace gmx
 120
 121 enum class GpuApiCallBehavior;
 122
 123 /* TODO error checking needs to be rewritten. We have 2 types of error checks needed
 124    based on where they occur in the code:
 125    - non performance-critical: these errors are unsafe to be ignored and must be
 126      _always_ checked for, e.g. initializations
 127    - performance critical: handling errors might hurt performance so care need to be taken
 128      when/if we should check for them at all, e.g. in cu_upload_X. However, we should be
 129      able to turn the check for these errors on!
 130
 131    Probably we'll need two sets of the macros below...
 132
 133  */
 134 #define CHECK_CUDA_ERRORS
 135
 136 #ifdef CHECK_CUDA_ERRORS
 137
 138 /*! Check for CUDA error on the return status of a CUDA RT API call. */
 139 #    define CU_RET_ERR(deviceError, msg)                                                            \
 140         do                                                                                          \
 141         {                                                                                           \
 142             if ((deviceError) != cudaSuccess)                                                       \
 143             {                                                                                       \
 144                 gmx_fatal(FARGS, "%s\n", ((msg) + gmx::getDeviceErrorString(deviceError)).c_str()); \
 145             }                                                                                       \
 146         } while (0)
 147
 148 #else /* CHECK_CUDA_ERRORS */
 149
 150 #    define CU_RET_ERR(status, msg) \
 151         do                          \
 152         {                           \
 153         } while (0)
 154
 155 #endif /* CHECK_CUDA_ERRORS */
 156
 157 // TODO: the 2 functions below are pretty much a constructor/destructor of a simple
 158 // GPU table object. There is also almost self-contained fetchFromParamLookupTable()
 159 // in cuda_kernel_utils.cuh. They could all live in a separate class/struct file.
 160
 161 /*! \brief Add a triplets stored in a float3 to an rvec variable.
 162  *
 163  * \param[out]  a Rvec to increment
 164  * \param[in]   b Float triplet to increment with.
 165  */
 166 static inline void rvec_inc(rvec a, const float3 b)
 167 {
 168     rvec tmp = { b.x, b.y, b.z };
 169     rvec_inc(a, tmp);
 170 }
 171
 172 /*! \brief  Returns true if all tasks in \p s have completed.
 173  *
 174  *  \param[in] deviceStream CUDA stream to check.
 175  *
 176  *  \returns True if all tasks enqueued in the stream \p deviceStream (at the time of this call) have completed.
 177  */
 178 static inline bool haveStreamTasksCompleted(const DeviceStream& deviceStream)
 179 {
 180     cudaError_t stat = cudaStreamQuery(deviceStream.stream());
 181
 182     if (stat == cudaErrorNotReady)
 183     {
 184         // work is still in progress in the stream
 185         return false;
 186     }
 187
 188     GMX_ASSERT(stat != cudaErrorInvalidResourceHandle,
 189                ("Stream identifier not valid. " + gmx::getDeviceErrorString(stat)).c_str());
 190
 191     // cudaSuccess and cudaErrorNotReady are the expected return values
 192     CU_RET_ERR(stat, "Unexpected cudaStreamQuery failure. ");
 193
 194     GMX_ASSERT(stat == cudaSuccess,
 195                ("Values other than cudaSuccess should have been explicitly handled. "
 196                 + gmx::getDeviceErrorString(stat))
 197                        .c_str());
 198
 199     return true;
 200 }
 201
 202 /* Kernel launch helpers */
 203
 204 /*! \brief
 205  * A function for setting up a single CUDA kernel argument.
 206  * This is the tail of the compile-time recursive function below.
 207  * It has to be seen by the compiler first.
 208  *
 209  * \tparam        totalArgsCount  Number of the kernel arguments
 210  * \tparam        KernelPtr       Kernel function handle type
 211  * \param[in]     argIndex        Index of the current argument
 212  */
 213 template<size_t totalArgsCount, typename KernelPtr>
 214 void prepareGpuKernelArgument(KernelPtr /*kernel*/,
 215                               std::array<void*, totalArgsCount>* /* kernelArgsPtr */,
 216                               size_t gmx_used_in_debug argIndex)
 217 {
 218     GMX_ASSERT(argIndex == totalArgsCount, "Tail expansion");
 219 }
 220
 221 /*! \brief
 222  * Compile-time recursive function for setting up a single CUDA kernel argument.
 223  * This function copies a kernel argument pointer \p argPtr into \p kernelArgsPtr,
 224  * and calls itself on the next argument, eventually calling the tail function above.
 225  *
 226  * \tparam        CurrentArg      Type of the current argument
 227  * \tparam        RemainingArgs   Types of remaining arguments after the current one
 228  * \tparam        totalArgsCount  Number of the kernel arguments
 229  * \tparam        KernelPtr       Kernel function handle type
 230  * \param[in]     kernel          Kernel function handle
 231  * \param[in,out] kernelArgsPtr   Pointer to the argument array to be filled in
 232  * \param[in]     argIndex        Index of the current argument
 233  * \param[in]     argPtr          Pointer to the current argument
 234  * \param[in]     otherArgsPtrs   Pack of pointers to arguments remaining to process after the current one
 235  */
 236 template<typename CurrentArg, typename... RemainingArgs, size_t totalArgsCount, typename KernelPtr>
 237 void prepareGpuKernelArgument(KernelPtr                          kernel,
 238                               std::array<void*, totalArgsCount>* kernelArgsPtr,
 239                               size_t                             argIndex,
 240                               const CurrentArg*                  argPtr,
 241                               const RemainingArgs*... otherArgsPtrs)
 242 {
 243     (*kernelArgsPtr)[argIndex] = const_cast<void*>(static_cast<const void*>(argPtr));
 244     prepareGpuKernelArgument(kernel, kernelArgsPtr, argIndex + 1, otherArgsPtrs...);
 245 }
 246
 247 /*! \brief
 248  * A wrapper function for setting up all the CUDA kernel arguments.
 249  * Calls the recursive functions above.
 250  *
 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.
 256  */
 257 template<typename KernelPtr, typename... Args>
 258 std::array<void*, sizeof...(Args)> prepareGpuKernelArguments(KernelPtr kernel,
 259                                                              const KernelLaunchConfig& /*config */,
 260                                                              const Args*... argsPtrs)
 261 {
 262     std::array<void*, sizeof...(Args)> kernelArgs;
 263     prepareGpuKernelArgument(kernel, &kernelArgs, 0, argsPtrs...);
 264     return kernelArgs;
 265 }
 266
 267 /*! \brief Launches the CUDA kernel and handles the errors.
 268  *
 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()
 276  * \throws gmx::InternalError on kernel launch failure
 277  */
 278 template<typename... Args>
 279 void launchGpuKernel(void (*kernel)(Args...),
 280                      const KernelLaunchConfig& config,
 281                      const DeviceStream&       deviceStream,
 282                      CommandEvent* /*timingEvent */,
 283                      const char*                               kernelName,
 284                      const std::array<void*, sizeof...(Args)>& kernelArgs)
 285 {
 286     dim3 blockSize(config.blockSize[0], config.blockSize[1], config.blockSize[2]);
 287     dim3 gridSize(config.gridSize[0], config.gridSize[1], config.gridSize[2]);
 288     cudaLaunchKernel(reinterpret_cast<void*>(kernel),
 289                      gridSize,
 290                      blockSize,
 291                      const_cast<void**>(kernelArgs.data()),
 292                      config.sharedMemorySize,
 293                      deviceStream.stream());
 294
 295     gmx::ensureNoPendingDeviceError("GPU kernel (" + std::string(kernelName)
 296                                     + ") failed to launch.");
 297 }
 298
 299 #endif