+//! \brief Interface for SYCL kernel function objects.
+class ISyclKernelFunctor
+{
+public:
+ //! \brief Virtual destructor.
+ virtual ~ISyclKernelFunctor() = default;
+ /*! \brief Set the kernel argument number \p argIndex to \p arg.
+ *
+ * \param argIndex Index of the argument. Maximum allowed value depends
+ * on the specific concrete class implementing this interface.
+ * \param arg Pointer to the argument value.
+ *
+ * \note Valid values of \p argIndex and types of \p arg depend on the
+ * specific concrete class implementing this interface. Passing
+ * illegal values is undefined behavior.
+ * \note Similar to \c clSetKernelArg, it is not safe to call this
+ * function on the same kernel object from multiple host threads.
+ */
+ virtual void setArg(size_t argIndex, void* arg) = 0;
+ /*! \brief Launch the kernel.
+ *
+ * \param config Work-group configuration.
+ * \param deviceStream \c DeviceStream to use.
+ */
+ virtual cl::sycl::event launch(const KernelLaunchConfig& /*config*/,
+ const DeviceStream& /*deviceStream*/) = 0;
+};
+
+
+/*! \brief
+ * A function for setting up a single SYCL kernel argument.
+ * This is the tail of the compile-time recursive function below.
+ * It has to be seen by the compiler first.
+ *
+ * \param[in] kernel Kernel function handle
+ * \param[in] argIndex Index of the current argument
+ */
+void inline prepareGpuKernelArgument(ISyclKernelFunctor* /*kernel*/, size_t /*argIndex*/) {}
+
+/*! \brief
+ * Compile-time recursive function for setting up a single SYCL kernel argument.
+ * This function uses one kernel argument pointer \p argPtr to call
+ * \c ISyclKernelFunctor::setArg, 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] 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(ISyclKernelFunctor* kernel,
+ size_t argIndex,
+ const CurrentArg* argPtr,
+ const RemainingArgs*... otherArgsPtrs)
+{
+ kernel->setArg(argIndex, const_cast<void*>(reinterpret_cast<const void*>(argPtr)));
+ prepareGpuKernelArgument(kernel, argIndex + 1, otherArgsPtrs...);
+}