2 * This file is part of the GROMACS molecular simulation package.
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.
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.
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.
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.
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.
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.
35 /*! \libinternal \file
36 * \brief Declare utility routines for SYCL
38 * \author Andrey Alekseenko <al42and@gmail.com>
41 #ifndef GMX_GPU_UTILS_SYCLUTILS_H
42 #define GMX_GPU_UTILS_SYCLUTILS_H
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"
53 enum class GpuApiCallBehavior;
56 * \brief SYCL GPU runtime data
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).
64 struct gmx_device_runtime_data_t
70 //! \brief Interface for SYCL kernel function objects.
71 class ISyclKernelFunctor
74 //! \brief Virtual destructor.
75 virtual ~ISyclKernelFunctor() = default;
76 /*! \brief Set the kernel argument number \p argIndex to \p arg.
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.
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.
88 virtual void setArg(size_t argIndex, void* arg) = 0;
89 /*! \brief Launch the kernel.
91 * \param config Work-group configuration.
92 * \param deviceStream \c DeviceStream to use.
94 virtual cl::sycl::event launch(const KernelLaunchConfig& /*config*/,
95 const DeviceStream& /*deviceStream*/) = 0;
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.
104 * \param[in] kernel Kernel function handle
105 * \param[in] argIndex Index of the current argument
107 void inline prepareGpuKernelArgument(ISyclKernelFunctor* /*kernel*/, size_t /*argIndex*/) {}
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.
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
122 template<typename CurrentArg, typename... RemainingArgs>
123 void prepareGpuKernelArgument(ISyclKernelFunctor* kernel,
125 const CurrentArg* argPtr,
126 const RemainingArgs*... otherArgsPtrs)
128 kernel->setArg(argIndex, const_cast<void*>(reinterpret_cast<const void*>(argPtr)));
129 prepareGpuKernelArgument(kernel, argIndex + 1, otherArgsPtrs...);
133 * A wrapper function for setting up all the SYCL kernel arguments.
134 * Calls the recursive functions above.
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.
142 template<typename... Args>
143 void* prepareGpuKernelArguments(void* kernel, const KernelLaunchConfig& /*config*/, const Args*... argsPtrs)
145 auto* kernelFunctor = reinterpret_cast<ISyclKernelFunctor*>(kernel);
146 prepareGpuKernelArgument(kernelFunctor, 0, argsPtrs...);
150 /*! \brief Launches the SYCL kernel and handles the errors.
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
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*/)
167 auto* kernelFunctor = reinterpret_cast<ISyclKernelFunctor*>(kernel);
168 cl::sycl::event event = kernelFunctor->launch(config, deviceStream);
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"
176 /*! \brief Pretend to check a SYCL stream for unfinished work (dummy implementation).
178 * \returns Not implemented in SYCL.
180 static inline bool haveStreamTasksCompleted(const DeviceStream& /* deviceStream */)
182 GMX_THROW(gmx::NotImplementedError("Not implemented on SYCL yet"));
185 # pragma clang diagnostic pop