prepareGpuKernelArguments() and launchGpuKernel() are added

[alexxy/gromacs.git] / src / gromacs / gpu_utils / oclutils.h

diff --git a/src/gromacs/gpu_utils/oclutils.h b/src/gromacs/gpu_utils/oclutils.h
index afe32388754d342c7b096dba517dc77cc58aba3f..ef7de9aa54dc60ac84e24ab651cdd0ef83d748b3 100644 (file)
--- a/src/gromacs/gpu_utils/oclutils.h
+++ b/src/gromacs/gpu_utils/oclutils.h
@@ -44,6 +44,8 @@
 #include <string>
 
 #include "gromacs/gpu_utils/gmxopencl.h"
+#include "gromacs/gpu_utils/gputraits_ocl.h"
+#include "gromacs/utility/exceptions.h"
 #include "gromacs/utility/gmxassert.h"
 
 enum class GpuApiCallBehavior;
@@ -177,4 +179,107 @@ static inline bool haveStreamTasksCompleted(cl_command_queue gmx_unused s)
     return false;
 }
 
+/* Kernel launch helpers */
+
+/*! \brief
+ * A function for setting up a single OpenCL kernel argument.
+ * This is the tail of the compile-time recursive function below.
+ * It has to be seen by the compiler first.
+ * As NB kernels might be using dynamic local memory as the last argument,
+ * this function also manages that, using sharedMemorySize from \p config.
+ *
+ * \param[in]     kernel          Kernel function handle
+ * \param[in]     config          Kernel configuration for launching
+ * \param[in]     argIndex        Index of the current argument
+ */
+void inline prepareGpuKernelArgument(cl_kernel                 kernel,
+                                     const KernelLaunchConfig &config,
+                                     size_t                    argIndex)
+{
+    if (config.sharedMemorySize > 0)
+    {
+        cl_int gmx_used_in_debug clError = clSetKernelArg(kernel, argIndex, config.sharedMemorySize, nullptr);
+        GMX_ASSERT(CL_SUCCESS == clError, ocl_get_error_string(clError).c_str());
+    }
+}
+
+/*! \brief
+ * Compile-time recursive function for setting up a single OpenCL kernel argument.
+ * This function uses one kernel argument pointer \p argPtr to call clSetKernelArg(),
+ * and calls itself on the next argument, eventually calling the tail function above.
+ *
+ * \tparam        CurrentArg      Type of the current argument
+ * \tparam        RemainingArgs   Types of remaining arguments after the current one
+ * \param[in]     kernel          Kernel function handle
+ * \param[in]     config          Kernel configuration for launching
+ * \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 CurrentArg, typename ... RemainingArgs>
+void prepareGpuKernelArgument(cl_kernel                 kernel,
+                              const KernelLaunchConfig &config,
+                              size_t                    argIndex,
+                              const CurrentArg         *argPtr,
+                              const RemainingArgs *...  otherArgsPtrs)
+{
+    cl_int gmx_used_in_debug clError = clSetKernelArg(kernel, argIndex, sizeof(CurrentArg), argPtr);
+    GMX_ASSERT(CL_SUCCESS == clError, ocl_get_error_string(clError).c_str());
+
+    prepareGpuKernelArgument(kernel, config, argIndex + 1, otherArgsPtrs ...);
+}
+
+/*! \brief
+ * A wrapper function for setting up all the OpenCL kernel arguments.
+ * Calls the recursive functions above.
+ *
+ * \tparam    Args            Types of all the kernel arguments
+ * \param[in] kernel          Kernel function handle
+ * \param[in] config          Kernel configuration for launching
+ * \param[in] argsPtrs        Pointers to all the kernel arguments
+ * \returns A handle for the prepared parameter pack to be used with launchGpuKernel() as the last argument
+ * - currently always nullptr for OpenCL, as it manages kernel/arguments association by itself.
+ */
+template <typename ... Args>
+void *prepareGpuKernelArguments(cl_kernel                 kernel,
+                                const KernelLaunchConfig &config,
+                                const Args *...           argsPtrs)
+{
+    prepareGpuKernelArgument(kernel, config, 0, argsPtrs ...);
+    return nullptr;
+}
+
+/*! \brief Launches the OpenCL kernel and handles the errors.
+ *
+ * \param[in] kernel          Kernel function handle
+ * \param[in] config          Kernel configuration for launching
+ * \param[in] timingEvent     Timing event, fetched from GpuRegionTimer
+ * \param[in] kernelName      Human readable kernel description, for error handling only
+ * \throws gmx::InternalError on kernel launch failure
+ */
+inline void launchGpuKernel(cl_kernel                 kernel,
+                            const KernelLaunchConfig &config,
+                            CommandEvent             *timingEvent,
+                            const char               *kernelName,
+                            const void                * /*kernelArgs*/)
+{
+    const int       workDimensions    = 3;
+    const size_t   *globalWorkOffset  = nullptr;
+    const size_t    waitListSize      = 0;
+    const cl_event *waitList          = nullptr;
+    size_t          globalWorkSize[3];
+    for (int i = 0; i < workDimensions; i++)
+    {
+        globalWorkSize[i] = config.gridSize[i] * config.blockSize[i];
+    }
+    cl_int clError = clEnqueueNDRangeKernel(config.stream, kernel, workDimensions, globalWorkOffset,
+                                            globalWorkSize, config.blockSize, waitListSize, waitList, timingEvent);
+    if (CL_SUCCESS != clError)
+    {
+        const std::string errorMessage = "GPU kernel (" +  std::string(kernelName) +
+            ") failed to launch: " + ocl_get_error_string(clError);
+        GMX_THROW(gmx::InternalError(errorMessage));
+    }
+}
+
 #endif