src/gromacs/gpu_utils/oclutils.h

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2014,2015,2016,2017,2018,2019, by the GROMACS development team, led by
   5  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   6  * and including many others, as listed in the AUTHORS file in the
   7  * top-level source directory and at http://www.gromacs.org.
   8  *
   9  * GROMACS is free software; you can redistribute it and/or
  10  * modify it under the terms of the GNU Lesser General Public License
  11  * as published by the Free Software Foundation; either version 2.1
  12  * of the License, or (at your option) any later version.
  13  *
  14  * GROMACS is distributed in the hope that it will be useful,
  15  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  16  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  17  * Lesser General Public License for more details.
  18  *
  19  * You should have received a copy of the GNU Lesser General Public
  20  * License along with GROMACS; if not, see
  21  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  22  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  23  *
  24  * If you want to redistribute modifications to GROMACS, please
  25  * consider that scientific software is very special. Version
  26  * control is crucial - bugs must be traceable. We will be happy to
  27  * consider code for inclusion in the official distribution, but
  28  * derived work must not be called official GROMACS. Details are found
  29  * in the README & COPYING files - if they are missing, get the
  30  * official version at http://www.gromacs.org.
  31  *
  32  * To help us fund GROMACS development, we humbly ask that you cite
  33  * the research papers on the package. Check out http://www.gromacs.org.
  34  */
  35 /*! \libinternal \file
  36  *  \brief Declare utility routines for OpenCL
  37  *
  38  *  \author Anca Hamuraru <anca@streamcomputing.eu>
  39  *  \inlibraryapi
  40  */
  41 #ifndef GMX_GPU_UTILS_OCLUTILS_H
  42 #define GMX_GPU_UTILS_OCLUTILS_H
  43
  44 #include <string>
  45
  46 #include "gromacs/gpu_utils/gmxopencl.h"
  47 #include "gromacs/gpu_utils/gputraits_ocl.h"
  48 #include "gromacs/utility/exceptions.h"
  49 #include "gromacs/utility/gmxassert.h"
  50
  51 enum class GpuApiCallBehavior;
  52
  53 /*! \brief OpenCL vendor IDs */
  54 typedef enum
  55 {
  56     OCL_VENDOR_NVIDIA = 0,
  57     OCL_VENDOR_AMD,
  58     OCL_VENDOR_INTEL,
  59     OCL_VENDOR_UNKNOWN
  60 } ocl_vendor_id_t;
  61
  62 /*! \internal
  63  * \brief OpenCL GPU device identificator
  64  *
  65  * An OpenCL device is identified by its ID.
  66  * The platform ID is also included for caching reasons.
  67  */
  68 typedef struct
  69 {
  70     cl_platform_id ocl_platform_id; /**< Platform ID */
  71     cl_device_id   ocl_device_id;   /**< Device ID */
  72 } ocl_gpu_id_t;
  73
  74 /*! \internal
  75  * \brief OpenCL device information.
  76  *
  77  * The OpenCL device information is queried and set at detection and contains
  78  * both information about the device/hardware returned by the runtime as well
  79  * as additional data like support status.
  80  */
  81 struct gmx_device_info_t
  82 {
  83     ocl_gpu_id_t    ocl_gpu_id;          /**< device ID assigned at detection   */
  84     char            device_name[256];    /**< device name */
  85     char            device_version[256]; /**< device version */
  86     char            device_vendor[256];  /**< device vendor */
  87     int             compute_units;       /**< number of compute units */
  88     int             adress_bits;         /**< number of adress bits the device is capable of */
  89     int             stat;                /**< device status takes values of e_gpu_detect_res_t */
  90     ocl_vendor_id_t vendor_e;            /**< device vendor as defined by ocl_vendor_id_t */
  91     size_t maxWorkItemSizes[3]; /**< workgroup size limits (CL_DEVICE_MAX_WORK_ITEM_SIZES) */
  92     size_t maxWorkGroupSize;    /**< workgroup total size limit (CL_DEVICE_MAX_WORK_GROUP_SIZE) */
  93 };
  94
  95 /*! \internal
  96  * \brief OpenCL GPU runtime data
  97  *
  98  * The device runtime data is meant to hold objects associated with a GROMACS rank's
  99  * (thread or process) use of a single device (multiple devices per rank is not
 100  * implemented). These objects should be constructed at ther point where a device
 101  * dets assigned to a rank and released at when this assignment is no longer valid
 102  * (i.e. at cleanup in the current implementation).
 103  *
 104  */
 105 struct gmx_device_runtime_data_t
 106 {
 107     cl_context context; /**< OpenCL context */
 108     cl_program program; /**< OpenCL program */
 109 };
 110
 111 /*! \brief Launches synchronous or asynchronous device to host memory copy.
 112  *
 113  *  If copy_event is not NULL, on return it will contain an event object
 114  *  identifying this particular device to host operation. The event can further
 115  *  be used to queue a wait for this operation or to query profiling information.
 116  */
 117 int ocl_copy_D2H(void*              h_dest,
 118                  cl_mem             d_src,
 119                  size_t             offset,
 120                  size_t             bytes,
 121                  GpuApiCallBehavior transferKind,
 122                  cl_command_queue   command_queue,
 123                  cl_event*          copy_event);
 124
 125
 126 /*! \brief Launches asynchronous device to host memory copy. */
 127 int ocl_copy_D2H_async(void*            h_dest,
 128                        cl_mem           d_src,
 129                        size_t           offset,
 130                        size_t           bytes,
 131                        cl_command_queue command_queue,
 132                        cl_event*        copy_event);
 133
 134 /*! \brief Launches synchronous or asynchronous host to device memory copy.
 135  *
 136  *  If copy_event is not NULL, on return it will contain an event object
 137  *  identifying this particular host to device operation. The event can further
 138  *  be used to queue a wait for this operation or to query profiling information.
 139  */
 140 int ocl_copy_H2D(cl_mem             d_dest,
 141                  const void*        h_src,
 142                  size_t             offset,
 143                  size_t             bytes,
 144                  GpuApiCallBehavior transferKind,
 145                  cl_command_queue   command_queue,
 146                  cl_event*          copy_event);
 147
 148 /*! \brief Launches asynchronous host to device memory copy. */
 149 int ocl_copy_H2D_async(cl_mem           d_dest,
 150                        const void*      h_src,
 151                        size_t           offset,
 152                        size_t           bytes,
 153                        cl_command_queue command_queue,
 154                        cl_event*        copy_event);
 155
 156 /*! \brief Launches synchronous host to device memory copy. */
 157 int ocl_copy_H2D_sync(cl_mem d_dest, const void* h_src, size_t offset, size_t bytes, cl_command_queue command_queue);
 158
 159 /*! \brief Allocate host memory in malloc style */
 160 void pmalloc(void** h_ptr, size_t nbytes);
 161
 162 /*! \brief Free host memory in malloc style */
 163 void pfree(void* h_ptr);
 164
 165 /*! \brief Convert error code to diagnostic string */
 166 std::string ocl_get_error_string(cl_int error);
 167
 168 /*! \brief Calls clFinish() in the stream \p s.
 169  *
 170  * \param[in] s stream to synchronize with
 171  */
 172 static inline void gpuStreamSynchronize(cl_command_queue s)
 173 {
 174     cl_int cl_error = clFinish(s);
 175     GMX_RELEASE_ASSERT(CL_SUCCESS == cl_error,
 176                        ("Error caught during clFinish:" + ocl_get_error_string(cl_error)).c_str());
 177 }
 178
 179 //! A debug checker to track cl_events being released correctly
 180 inline void ensureReferenceCount(const cl_event& event, unsigned int refCount)
 181 {
 182 #ifndef NDEBUG
 183     cl_int clError = clGetEventInfo(event, CL_EVENT_REFERENCE_COUNT, sizeof(refCount), &refCount, nullptr);
 184     GMX_ASSERT(CL_SUCCESS == clError, ocl_get_error_string(clError).c_str());
 185     GMX_ASSERT(refCount == refCount, "Unexpected reference count");
 186 #else
 187     GMX_UNUSED_VALUE(event);
 188     GMX_UNUSED_VALUE(refCount);
 189 #endif
 190 }
 191
 192 /*! \brief Pretend to synchronize an OpenCL stream (dummy implementation).
 193  *
 194  * \param[in] s queue to check
 195  *
 196  *  \returns     True if all tasks enqueued in the stream \p s (at the time of this call) have completed.
 197  */
 198 static inline bool haveStreamTasksCompleted(cl_command_queue gmx_unused s)
 199 {
 200     GMX_RELEASE_ASSERT(false, "haveStreamTasksCompleted is not implemented for OpenCL");
 201     return false;
 202 }
 203
 204 /* Kernel launch helpers */
 205
 206 /*! \brief
 207  * A function for setting up a single OpenCL kernel argument.
 208  * This is the tail of the compile-time recursive function below.
 209  * It has to be seen by the compiler first.
 210  * As NB kernels might be using dynamic local memory as the last argument,
 211  * this function also manages that, using sharedMemorySize from \p config.
 212  *
 213  * \param[in]     kernel          Kernel function handle
 214  * \param[in]     config          Kernel configuration for launching
 215  * \param[in]     argIndex        Index of the current argument
 216  */
 217 void inline prepareGpuKernelArgument(cl_kernel kernel, const KernelLaunchConfig& config, size_t argIndex)
 218 {
 219     if (config.sharedMemorySize > 0)
 220     {
 221         cl_int gmx_used_in_debug clError =
 222                 clSetKernelArg(kernel, argIndex, config.sharedMemorySize, nullptr);
 223         GMX_ASSERT(CL_SUCCESS == clError, ocl_get_error_string(clError).c_str());
 224     }
 225 }
 226
 227 /*! \brief
 228  * Compile-time recursive function for setting up a single OpenCL kernel argument.
 229  * This function uses one kernel argument pointer \p argPtr to call clSetKernelArg(),
 230  * and calls itself on the next argument, eventually calling the tail function above.
 231  *
 232  * \tparam        CurrentArg      Type of the current argument
 233  * \tparam        RemainingArgs   Types of remaining arguments after the current one
 234  * \param[in]     kernel          Kernel function handle
 235  * \param[in]     config          Kernel configuration for launching
 236  * \param[in]     argIndex        Index of the current argument
 237  * \param[in]     argPtr          Pointer to the current argument
 238  * \param[in]     otherArgsPtrs   Pack of pointers to arguments remaining to process after the current one
 239  */
 240 template<typename CurrentArg, typename... RemainingArgs>
 241 void prepareGpuKernelArgument(cl_kernel                 kernel,
 242                               const KernelLaunchConfig& config,
 243                               size_t                    argIndex,
 244                               const CurrentArg*         argPtr,
 245                               const RemainingArgs*... otherArgsPtrs)
 246 {
 247     cl_int gmx_used_in_debug clError = clSetKernelArg(kernel, argIndex, sizeof(CurrentArg), argPtr);
 248     GMX_ASSERT(CL_SUCCESS == clError, ocl_get_error_string(clError).c_str());
 249
 250     // Assert on types not allowed to be passed to a kernel
 251     // (as per section 6.9 of the OpenCL spec).
 252     static_assert(!std::is_same<CurrentArg, bool>::value && !std::is_same<CurrentArg, size_t>::value
 253                           && !std::is_same<CurrentArg, ptrdiff_t>::value
 254                           && !std::is_same<CurrentArg, intptr_t>::value
 255                           && !std::is_same<CurrentArg, uintptr_t>::value,
 256                   "Invalid type passed to OpenCL kernel functions (see OpenCL spec section 6.9).");
 257
 258     prepareGpuKernelArgument(kernel, config, argIndex + 1, otherArgsPtrs...);
 259 }
 260
 261 /*! \brief
 262  * A wrapper function for setting up all the OpenCL kernel arguments.
 263  * Calls the recursive functions above.
 264  *
 265  * \tparam    Args            Types of all the kernel arguments
 266  * \param[in] kernel          Kernel function handle
 267  * \param[in] config          Kernel configuration for launching
 268  * \param[in] argsPtrs        Pointers to all the kernel arguments
 269  * \returns A handle for the prepared parameter pack to be used with launchGpuKernel() as the last argument
 270  * - currently always nullptr for OpenCL, as it manages kernel/arguments association by itself.
 271  */
 272 template<typename... Args>
 273 void* prepareGpuKernelArguments(cl_kernel kernel, const KernelLaunchConfig& config, const Args*... argsPtrs)
 274 {
 275     prepareGpuKernelArgument(kernel, config, 0, argsPtrs...);
 276     return nullptr;
 277 }
 278
 279 /*! \brief Launches the OpenCL kernel and handles the errors.
 280  *
 281  * \param[in] kernel          Kernel function handle
 282  * \param[in] config          Kernel configuration for launching
 283  * \param[in] timingEvent     Timing event, fetched from GpuRegionTimer
 284  * \param[in] kernelName      Human readable kernel description, for error handling only
 285  * \throws gmx::InternalError on kernel launch failure
 286  */
 287 inline void launchGpuKernel(cl_kernel                 kernel,
 288                             const KernelLaunchConfig& config,
 289                             CommandEvent*             timingEvent,
 290                             const char*               kernelName,
 291                             const void* /*kernelArgs*/)
 292 {
 293     const int       workDimensions   = 3;
 294     const size_t*   globalWorkOffset = nullptr;
 295     const size_t    waitListSize     = 0;
 296     const cl_event* waitList         = nullptr;
 297     size_t          globalWorkSize[3];
 298     for (int i = 0; i < workDimensions; i++)
 299     {
 300         globalWorkSize[i] = config.gridSize[i] * config.blockSize[i];
 301     }
 302     cl_int clError = clEnqueueNDRangeKernel(config.stream, kernel, workDimensions, globalWorkOffset,
 303                                             globalWorkSize, config.blockSize, waitListSize,
 304                                             waitList, timingEvent);
 305     if (CL_SUCCESS != clError)
 306     {
 307         const std::string errorMessage = "GPU kernel (" + std::string(kernelName)
 308                                          + ") failed to launch: " + ocl_get_error_string(clError);
 309         GMX_THROW(gmx::InternalError(errorMessage));
 310     }
 311 }
 312
 313 #endif