/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2012,2014,2015,2016,2017, by the GROMACS development team, led by
+ * Copyright (c) 2012,2014,2015,2016,2017 by the GROMACS development team.
+ * Copyright (c) 2018,2019,2020,2021, by the GROMACS development team, led by
* Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
* and including many others, as listed in the AUTHORS file in the
* top-level source directory and at http://www.gromacs.org.
#ifndef GMX_GPU_UTILS_CUDAUTILS_CUH
#define GMX_GPU_UTILS_CUDAUTILS_CUH
-#include "config.h"
-
#include <stdio.h>
-#if HAVE_NVML
-#include <nvml.h>
-#endif /* HAVE_NVML */
+#include <array>
#include <string>
+#include <type_traits>
+#include "gromacs/gpu_utils/device_stream.h"
+#include "gromacs/gpu_utils/gputraits.cuh"
#include "gromacs/math/vec.h"
#include "gromacs/math/vectypes.h"
+#include "gromacs/utility/exceptions.h"
#include "gromacs/utility/fatalerror.h"
#include "gromacs/utility/gmxassert.h"
#include "gromacs/utility/stringutil.h"
namespace
{
+/*! \brief Add the API information on the specific error to the error message.
+ *
+ * \param[in] deviceError The error to assert cudaSuccess on.
+ *
+ * \returns A description of the API error. Returns '(CUDA error #0 (cudaSuccess): no error)' in case deviceError is cudaSuccess.
+ */
+inline std::string getDeviceErrorString(const cudaError_t deviceError)
+{
+ return formatString("CUDA error #%d (%s): %s.",
+ deviceError,
+ cudaGetErrorName(deviceError),
+ cudaGetErrorString(deviceError));
+}
+
+/*! \brief Check if API returned an error and throw an exception with information on it.
+ *
+ * \param[in] deviceError The error to assert cudaSuccess on.
+ * \param[in] errorMessage Undecorated error message.
+ *
+ * \throws InternalError if deviceError is not a success.
+ */
+inline void checkDeviceError(const cudaError_t deviceError, const std::string& errorMessage)
+{
+ if (deviceError != cudaSuccess)
+ {
+ GMX_THROW(gmx::InternalError(errorMessage + " " + getDeviceErrorString(deviceError)));
+ }
+}
+
/*! \brief Helper function to ensure no pending error silently
* disrupts error handling.
*
* Asserts in a debug build if an unhandled error is present. Issues a
* warning at run time otherwise.
*
- * \todo This is similar to CU_CHECK_PREV_ERR, which should be
- * consolidated.
+ * \param[in] errorMessage Undecorated error message.
*/
-static inline void ensureNoPendingCudaError(const char *errorMessage)
+inline void ensureNoPendingDeviceError(const std::string& errorMessage)
{
// Ensure there is no pending error that would otherwise affect
// the behaviour of future error handling.
- cudaError_t stat = cudaGetLastError();
- if (stat == cudaSuccess)
+ cudaError_t deviceError = cudaGetLastError();
+ if (deviceError == cudaSuccess)
{
return;
}
// If we would find an error in a release build, we do not know
// what is appropriate to do about it, so assert only for debug
// builds.
- auto fullMessage = formatString("%s An unhandled error from a previous CUDA operation was detected. %s: %s",
- errorMessage, cudaGetErrorName(stat), cudaGetErrorString(stat));
- GMX_ASSERT(stat == cudaSuccess, fullMessage.c_str());
+ const std::string fullErrorMessage =
+ errorMessage + " An unhandled error from a previous CUDA operation was detected. "
+ + gmx::getDeviceErrorString(deviceError);
+ GMX_ASSERT(deviceError == cudaSuccess, fullErrorMessage.c_str());
// TODO When we evolve a better logging framework, use that
// for release-build error reporting.
- gmx_warning(fullMessage.c_str());
+ gmx_warning("%s", fullErrorMessage.c_str());
}
-} // namespace
-} // namespace
+} // namespace
+} // namespace gmx
enum class GpuApiCallBehavior;
#ifdef CHECK_CUDA_ERRORS
/*! Check for CUDA error on the return status of a CUDA RT API call. */
-#define CU_RET_ERR(status, msg) \
- do { \
- if (status != cudaSuccess) \
- { \
- gmx_fatal(FARGS, "%s: %s\n", msg, cudaGetErrorString(status)); \
- } \
- } while (0)
-
-/*! Check for any previously occurred uncaught CUDA error. */
-#define CU_CHECK_PREV_ERR() \
- do { \
- cudaError_t _CU_CHECK_PREV_ERR_status = cudaGetLastError(); \
- if (_CU_CHECK_PREV_ERR_status != cudaSuccess) { \
- gmx_warning("Just caught a previously occurred CUDA error (%s), will try to continue.", cudaGetErrorString(_CU_CHECK_PREV_ERR_status)); \
- } \
- } while (0)
-
-/*! Check for any previously occurred uncaught CUDA error
- -- aimed at use after kernel calls. */
-#define CU_LAUNCH_ERR(msg) \
- do { \
- cudaError_t _CU_LAUNCH_ERR_status = cudaGetLastError(); \
- if (_CU_LAUNCH_ERR_status != cudaSuccess) { \
- gmx_fatal(FARGS, "Error while launching kernel %s: %s\n", msg, cudaGetErrorString(_CU_LAUNCH_ERR_status)); \
- } \
- } while (0)
-
-/*! Synchronize with GPU and check for any previously occurred uncaught CUDA error
- -- aimed at use after kernel calls. */
-#define CU_LAUNCH_ERR_SYNC(msg) \
- do { \
- cudaError_t _CU_SYNC_LAUNCH_ERR_status = cudaThreadSynchronize(); \
- if (_CU_SYNC_LAUNCH_ERR_status != cudaSuccess) { \
- gmx_fatal(FARGS, "Error while launching kernel %s: %s\n", msg, cudaGetErrorString(_CU_SYNC_LAUNCH_ERR_status)); \
- } \
- } while (0)
+# define CU_RET_ERR(deviceError, msg) \
+ do \
+ { \
+ if ((deviceError) != cudaSuccess) \
+ { \
+ gmx_fatal(FARGS, "%s\n", ((msg) + gmx::getDeviceErrorString(deviceError)).c_str()); \
+ } \
+ } while (0)
#else /* CHECK_CUDA_ERRORS */
-#define CU_RET_ERR(status, msg) do { } while (0)
-#define CU_CHECK_PREV_ERR() do { } while (0)
-#define CU_LAUNCH_ERR(msg) do { } while (0)
-#define CU_LAUNCH_ERR_SYNC(msg) do { } while (0)
-#define HANDLE_NVML_RET_ERR(status, msg) do { } while (0)
+# define CU_RET_ERR(status, msg) \
+ do \
+ { \
+ } while (0)
#endif /* CHECK_CUDA_ERRORS */
-/*! \brief CUDA device information.
- *
- * The CUDA device information is queried and set at detection and contains
- * both information about the device/hardware returned by the runtime as well
- * as additional data like support status.
+// TODO: the 2 functions below are pretty much a constructor/destructor of a simple
+// GPU table object. There is also almost self-contained fetchFromParamLookupTable()
+// in cuda_kernel_utils.cuh. They could all live in a separate class/struct file.
+
+/*! \brief Add a triplets stored in a float3 to an rvec variable.
*
- * \todo extract an object to manage NVML details
+ * \param[out] a Rvec to increment
+ * \param[in] b Float triplet to increment with.
*/
-struct gmx_device_info_t
+static inline void rvec_inc(rvec a, const float3 b)
{
- int id; /* id of the CUDA device */
- cudaDeviceProp prop; /* CUDA device properties */
- int stat; /* result of the device check */
- unsigned int nvml_orig_app_sm_clock; /* The original SM clock before we changed it */
- unsigned int nvml_orig_app_mem_clock; /* The original memory clock before we changed it */
- gmx_bool nvml_app_clocks_changed; /* If application clocks have been changed */
- unsigned int nvml_set_app_sm_clock; /* The SM clock we set */
- unsigned int nvml_set_app_mem_clock; /* The memory clock we set */
-#if HAVE_NVML
- nvmlDevice_t nvml_device_id; /* NVML device id */
- // TODO This can become a bool with a more useful name
- nvmlEnableState_t nvml_is_restricted; /* Status of application clocks permission */
-#endif /* HAVE_NVML */
-};
-
-/*! Launches synchronous or asynchronous device to host memory copy.
- *
- * The copy is launched in stream s or if not specified, in stream 0.
- */
-int cu_copy_D2H(void *h_dest, void *d_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t s /*= 0*/);
-
-/*! Launches synchronous host to device memory copy in stream 0. */
-int cu_copy_D2H_sync(void * /*h_dest*/, void * /*d_src*/, size_t /*bytes*/);
-
-/*! Launches asynchronous host to device memory copy in stream s. */
-int cu_copy_D2H_async(void * /*h_dest*/, void * /*d_src*/, size_t /*bytes*/, cudaStream_t /*s = 0*/);
+ rvec tmp = { b.x, b.y, b.z };
+ rvec_inc(a, tmp);
+}
-/*! Launches synchronous or asynchronous host to device memory copy.
+/*! \brief Returns true if all tasks in \p s have completed.
+ *
+ * \param[in] deviceStream CUDA stream to check.
*
- * The copy is launched in stream s or if not specified, in stream 0.
+ * \returns True if all tasks enqueued in the stream \p deviceStream (at the time of this call) have completed.
*/
-int cu_copy_H2D(void *d_dest, void *h_src, size_t bytes, GpuApiCallBehavior transferKind, cudaStream_t /*s = 0*/);
+static inline bool haveStreamTasksCompleted(const DeviceStream& deviceStream)
+{
+ cudaError_t stat = cudaStreamQuery(deviceStream.stream());
-/*! Launches synchronous host to device memory copy. */
-int cu_copy_H2D_sync(void * /*d_dest*/, void * /*h_src*/, size_t /*bytes*/);
+ if (stat == cudaErrorNotReady)
+ {
+ // work is still in progress in the stream
+ return false;
+ }
-/*! Launches asynchronous host to device memory copy in stream s. */
-int cu_copy_H2D_async(void * /*d_dest*/, void * /*h_src*/, size_t /*bytes*/, cudaStream_t /*s = 0*/);
+ GMX_ASSERT(stat != cudaErrorInvalidResourceHandle,
+ ("Stream identifier not valid. " + gmx::getDeviceErrorString(stat)).c_str());
-/*! Frees device memory and resets the size and allocation size to -1. */
-void cu_free_buffered(void *d_ptr, int *n = NULL, int *nalloc = NULL);
+ // cudaSuccess and cudaErrorNotReady are the expected return values
+ CU_RET_ERR(stat, "Unexpected cudaStreamQuery failure. ");
-/*! Reallocates the device memory and copies data from the host. */
-void cu_realloc_buffered(void **d_dest, void *h_src,
- size_t type_size,
- int *curr_size, int *curr_alloc_size,
- int req_size,
- cudaStream_t s,
- bool bAsync);
+ GMX_ASSERT(stat == cudaSuccess,
+ ("Values other than cudaSuccess should have been explicitly handled. "
+ + gmx::getDeviceErrorString(stat))
+ .c_str());
-// TODO: the 2 functions below are pretty much a constructor/destructor of a simple
-// GPU table object. There is also almost self-contained fetchFromParamLookupTable()
-// in cuda_kernel_utils.cuh. They could all live in a separate class/struct file,
-// granted storing static texture references in there does not pose problems.
+ return true;
+}
-/*! \brief Initialize parameter lookup table.
- *
- * Initializes device memory, copies data from host and binds
- * a texture to allocated device memory to be used for parameter lookup.
+/* Kernel launch helpers */
+
+/*! \brief
+ * A function for setting up a single CUDA kernel argument.
+ * This is the tail of the compile-time recursive function below.
+ * It has to be seen by the compiler first.
*
- * \tparam[in] T Raw data type
- * \param[out] d_ptr device pointer to the memory to be allocated
- * \param[out] texObj texture object to be initialized
- * \param[out] texRef texture reference to be initialized
- * \param[in] h_ptr pointer to the host memory to be uploaded to the device
- * \param[in] numElem number of elements in the h_ptr
- * \param[in] devInfo pointer to the info struct of the device in use
+ * \tparam totalArgsCount Number of the kernel arguments
+ * \tparam KernelPtr Kernel function handle type
+ * \param[in] argIndex Index of the current argument
*/
-template <typename T>
-void initParamLookupTable(T * &d_ptr,
- cudaTextureObject_t &texObj,
- const struct texture<T, 1, cudaReadModeElementType> *texRef,
- const T *h_ptr,
- int numElem,
- const gmx_device_info_t *devInfo);
-
-/*! \brief Destroy parameter lookup table.
- *
- * Unbinds texture reference/object, deallocates device memory.
+template<size_t totalArgsCount, typename KernelPtr>
+void prepareGpuKernelArgument(KernelPtr /*kernel*/,
+ std::array<void*, totalArgsCount>* /* kernelArgsPtr */,
+ size_t gmx_used_in_debug argIndex)
+{
+ GMX_ASSERT(argIndex == totalArgsCount, "Tail expansion");
+}
+
+/*! \brief
+ * Compile-time recursive function for setting up a single CUDA kernel argument.
+ * This function copies a kernel argument pointer \p argPtr into \p kernelArgsPtr,
+ * and calls itself on the next argument, eventually calling the tail function above.
*
- * \tparam[in] T Raw data type
- * \param[in] d_ptr Device pointer to the memory to be deallocated
- * \param[in] texObj Texture object to be deinitialized
- * \param[in] texRef Texture reference to be deinitialized
- * \param[in] devInfo Pointer to the info struct of the device in use
+ * \tparam CurrentArg Type of the current argument
+ * \tparam RemainingArgs Types of remaining arguments after the current one
+ * \tparam totalArgsCount Number of the kernel arguments
+ * \tparam KernelPtr Kernel function handle type
+ * \param[in] kernel Kernel function handle
+ * \param[in,out] kernelArgsPtr Pointer to the argument array to be filled in
+ * \param[in] argIndex Index of the current argument
+ * \param[in] argPtr Pointer to the current argument
+ * \param[in] otherArgsPtrs Pack of pointers to arguments remaining to process after the current one
*/
-template <typename T>
-void destroyParamLookupTable(T *d_ptr,
- cudaTextureObject_t texObj,
- const struct texture<T, 1, cudaReadModeElementType> *texRef,
- const gmx_device_info_t *devInfo);
+template<typename CurrentArg, typename... RemainingArgs, size_t totalArgsCount, typename KernelPtr>
+void prepareGpuKernelArgument(KernelPtr kernel,
+ std::array<void*, totalArgsCount>* kernelArgsPtr,
+ size_t argIndex,
+ const CurrentArg* argPtr,
+ const RemainingArgs*... otherArgsPtrs)
+{
+ (*kernelArgsPtr)[argIndex] = const_cast<void*>(static_cast<const void*>(argPtr));
+ prepareGpuKernelArgument(kernel, kernelArgsPtr, argIndex + 1, otherArgsPtrs...);
+}
-/*! \brief Add a triplets stored in a float3 to an rvec variable.
+/*! \brief
+ * A wrapper function for setting up all the CUDA kernel arguments.
+ * Calls the recursive functions above.
*
- * \param[out] a Rvec to increment
- * \param[in] b Float triplet to increment with.
+ * \tparam KernelPtr Kernel function handle type
+ * \tparam Args Types of all the kernel arguments
+ * \param[in] kernel Kernel function handle
+ * \param[in] argsPtrs Pointers to all the kernel arguments
+ * \returns A prepared parameter pack to be used with launchGpuKernel() as the last argument.
*/
-static inline void rvec_inc(rvec a, const float3 b)
+template<typename KernelPtr, typename... Args>
+std::array<void*, sizeof...(Args)> prepareGpuKernelArguments(KernelPtr kernel,
+ const KernelLaunchConfig& /*config */,
+ const Args*... argsPtrs)
{
- rvec tmp = {b.x, b.y, b.z};
- rvec_inc(a, tmp);
+ std::array<void*, sizeof...(Args)> kernelArgs;
+ prepareGpuKernelArgument(kernel, &kernelArgs, 0, argsPtrs...);
+ return kernelArgs;
}
-/*! \brief Calls cudaStreamSynchronize() in the stream \p s.
+/*! \brief Launches the CUDA kernel and handles the errors.
*
- * \param[in] s stream to synchronize with
+ * \tparam Args Types of all the kernel arguments
+ * \param[in] kernel Kernel function handle
+ * \param[in] config Kernel configuration for launching
+ * \param[in] deviceStream GPU stream to launch kernel in
+ * \param[in] kernelName Human readable kernel description, for error handling only
+ * \param[in] kernelArgs Array of the pointers to the kernel arguments, prepared by
+ * prepareGpuKernelArguments()
+ * \throws gmx::InternalError on kernel launch failure
*/
-static inline void gpuStreamSynchronize(cudaStream_t s)
+template<typename... Args>
+void launchGpuKernel(void (*kernel)(Args...),
+ const KernelLaunchConfig& config,
+ const DeviceStream& deviceStream,
+ CommandEvent* /*timingEvent */,
+ const char* kernelName,
+ const std::array<void*, sizeof...(Args)>& kernelArgs)
{
- cudaError_t stat = cudaStreamSynchronize(s);
- CU_RET_ERR(stat, "cudaStreamSynchronize failed");
+ dim3 blockSize(config.blockSize[0], config.blockSize[1], config.blockSize[2]);
+ dim3 gridSize(config.gridSize[0], config.gridSize[1], config.gridSize[2]);
+ cudaLaunchKernel(reinterpret_cast<void*>(kernel),
+ gridSize,
+ blockSize,
+ const_cast<void**>(kernelArgs.data()),
+ config.sharedMemorySize,
+ deviceStream.stream());
+
+ gmx::ensureNoPendingDeviceError("GPU kernel (" + std::string(kernelName)
+ + ") failed to launch.");
}
#endif
biod.pnpi.spb.ru / alexxy / gromacs.git / blobdiff