src/gromacs/gpu_utils/syclutils.h

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2020,2021, 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 SYCL
  37  *
  38  *  \author Andrey Alekseenko <al42and@gmail.com>
  39  *  \inlibraryapi
  40  */
  41 #ifndef GMX_GPU_UTILS_SYCLUTILS_H
  42 #define GMX_GPU_UTILS_SYCLUTILS_H
  43
  44 #include <string>
  45
  46 #include "gromacs/gpu_utils/gmxsycl.h"
  47 #include "gromacs/gpu_utils/gputraits.h"
  48 #include "gromacs/utility/exceptions.h"
  49 #include "gromacs/utility/gmxassert.h"
  50 #include "gromacs/utility/stringutil.h"
  51
  52 class DeviceStream;
  53 enum class GpuApiCallBehavior;
  54
  55 /*! \internal
  56  * \brief SYCL GPU runtime data
  57  *
  58  * The device runtime data is meant to hold objects associated with a GROMACS rank's
  59  * (thread or process) use of a single device (multiple devices per rank is not
  60  * implemented). These objects should be constructed at the point where a device
  61  * gets assigned to a rank and released at when this assignment is no longer valid
  62  * (i.e. at cleanup in the current implementation).
  63  */
  64 struct gmx_device_runtime_data_t
  65 {
  66 };
  67
  68 #ifndef DOXYGEN
  69
  70 //! \brief Interface for SYCL kernel function objects.
  71 class ISyclKernelFunctor
  72 {
  73 public:
  74     //! \brief Virtual destructor.
  75     virtual ~ISyclKernelFunctor() = default;
  76     /*! \brief Set the kernel argument number \p argIndex to \p arg.
  77      *
  78      * \param argIndex Index of the argument. Maximum allowed value depends
  79      *                 on the specific concrete class implementing this interface.
  80      * \param arg      Pointer to the argument value.
  81      *
  82      * \note Valid values of \p argIndex and types of \p arg depend on the
  83      *       specific concrete class implementing this interface. Passing
  84      *       illegal values is undefined behavior.
  85      * \note Similar to \c clSetKernelArg, it is not safe to call this
  86      *       function on the same kernel object from multiple host threads.
  87      */
  88     virtual void setArg(size_t argIndex, void* arg) = 0;
  89     /*! \brief Launch the kernel.
  90      *
  91      * \param config       Work-group configuration.
  92      * \param deviceStream \c DeviceStream to use.
  93      */
  94     virtual cl::sycl::event launch(const KernelLaunchConfig& /*config*/,
  95                                    const DeviceStream& /*deviceStream*/) = 0;
  96 };
  97
  98
  99 /*! \brief
 100  * A function for setting up a single SYCL kernel argument.
 101  * This is the tail of the compile-time recursive function below.
 102  * It has to be seen by the compiler first.
 103  *
 104  * \param[in]     kernel          Kernel function handle
 105  * \param[in]     argIndex        Index of the current argument
 106  */
 107 void inline prepareGpuKernelArgument(ISyclKernelFunctor* /*kernel*/, size_t /*argIndex*/) {}
 108
 109 /*! \brief
 110  * Compile-time recursive function for setting up a single SYCL kernel argument.
 111  * This function uses one kernel argument pointer \p argPtr to call
 112  * \c ISyclKernelFunctor::setArg, and calls itself on the next argument, eventually
 113  * calling the tail function above.
 114  *
 115  * \tparam        CurrentArg      Type of the current argument
 116  * \tparam        RemainingArgs   Types of remaining arguments after the current one
 117  * \param[in]     kernel          Kernel function handle
 118  * \param[in]     argIndex        Index of the current argument
 119  * \param[in]     argPtr          Pointer to the current argument
 120  * \param[in]     otherArgsPtrs   Pack of pointers to arguments remaining to process after the current one
 121  */
 122 template<typename CurrentArg, typename... RemainingArgs>
 123 void prepareGpuKernelArgument(ISyclKernelFunctor* kernel,
 124                               size_t              argIndex,
 125                               const CurrentArg*   argPtr,
 126                               const RemainingArgs*... otherArgsPtrs)
 127 {
 128     kernel->setArg(argIndex, const_cast<void*>(reinterpret_cast<const void*>(argPtr)));
 129     prepareGpuKernelArgument(kernel, argIndex + 1, otherArgsPtrs...);
 130 }
 131
 132 /*! \brief
 133  * A wrapper function for setting up all the SYCL kernel arguments.
 134  * Calls the recursive functions above.
 135  *
 136  * \tparam    Args            Types of all the kernel arguments
 137  * \param[in] kernel          Kernel function handle
 138  * \param[in] config          Kernel configuration for launching
 139  * \param[in] argsPtrs        Pointers to all the kernel arguments
 140  * \returns A dummy value to be used with launchGpuKernel() as the last argument.
 141  */
 142 template<typename... Args>
 143 void* prepareGpuKernelArguments(void* kernel, const KernelLaunchConfig& /*config*/, const Args*... argsPtrs)
 144 {
 145     auto* kernelFunctor = reinterpret_cast<ISyclKernelFunctor*>(kernel);
 146     prepareGpuKernelArgument(kernelFunctor, 0, argsPtrs...);
 147     return nullptr;
 148 }
 149
 150 /*! \brief Launches the SYCL kernel and handles the errors.
 151  *
 152  * \param[in] kernel          Kernel function handle
 153  * \param[in] config          Kernel configuration for launching
 154  * \param[in] deviceStream    GPU stream to launch kernel in
 155  * \param[in] timingEvent     Timing event, fetched from GpuRegionTimer. Unused.
 156  * \param[in] kernelName      Human readable kernel description, for error handling only. Unused.
 157  * \param[in] kernelArgs      Unused.
 158  * \throws gmx::InternalError on kernel launch failure
 159  */
 160 inline void launchGpuKernel(void*                     kernel,
 161                             const KernelLaunchConfig& config,
 162                             const DeviceStream&       deviceStream,
 163                             CommandEvent* /*timingEvent*/,
 164                             const char* /*kernelName*/,
 165                             const void* /*kernelArgs*/)
 166 {
 167     auto*           kernelFunctor = reinterpret_cast<ISyclKernelFunctor*>(kernel);
 168     cl::sycl::event event         = kernelFunctor->launch(config, deviceStream);
 169 }
 170
 171 /* To properly mark function as [[noreturn]], we must do it everywhere it is declared, which
 172  * will pollute common headers.*/
 173 #    pragma clang diagnostic push
 174 #    pragma clang diagnostic ignored "-Wmissing-noreturn"
 175
 176 /*! \brief Pretend to check a SYCL stream for unfinished work (dummy implementation).
 177  *
 178  *  \returns  Not implemented in SYCL.
 179  */
 180 static inline bool haveStreamTasksCompleted(const DeviceStream& /* deviceStream */)
 181 {
 182     GMX_THROW(gmx::NotImplementedError("Not implemented on SYCL yet"));
 183 }
 184
 185 #    pragma clang diagnostic pop
 186
 187 #endif // !DOXYGEN
 188
 189 #endif