corecel/sys/KernelLauncher.device.hh

0001 //------------------------------ -*- C++ -*- -------------------------------//
0002 // Copyright Celeritas contributors: see top-level COPYRIGHT file for details
0003 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
0004 //---------------------------------------------------------------------------//
0005 //! \file corecel/sys/KernelLauncher.device.hh
0006 //---------------------------------------------------------------------------//
0007 #pragma once
0008
0009 #include <string>
0010
0011 #include "corecel/DeviceRuntimeApi.hh"
0012
0013 #include "corecel/Types.hh"
0014 #include "corecel/cont/Range.hh"
0015 #include "corecel/sys/ScopedProfiling.hh"
0016
0017 #include "Device.hh"
0018 #include "KernelParamCalculator.device.hh"
0019 #include "Stream.hh"
0020 #include "ThreadId.hh"
0021
0022 #include "detail/KernelLauncherImpl.device.hh"
0023
0024 namespace celeritas
0025 {
0026 //---------------------------------------------------------------------------//
0027 /*!
0028  * Profile and launch Celeritas kernels.
0029  *
0030  * The template argument \c F may define a member type named \c Applier.
0031  * \c F::Applier should have up to two static constexpr int variables named
0032  * \c max_block_size and/or \c min_warps_per_eu.
0033  * If present, the kernel will use appropriate \c __launch_bounds__.
0034  * If \c F::Applier::min_warps_per_eu exists then \c F::Applier::max_block_size
0035  * must also be present or we get a compile error.
0036  *
0037  * The semantics of the second \c __launch_bounds__ argument differs between
0038  * CUDA and HIP.  \c KernelLauncher expects HIP semantics. If Celeritas is
0039  * built targeting CUDA, it will automatically convert that argument to match
0040  * CUDA semantics.
0041  *
0042  * The CUDA-specific 3rd argument \c maxBlocksPerCluster is not supported.
0043  *
0044  * Example:
0045  * \code
0046  void launch_kernel(DeviceParams const& params, size_type count) const
0047  {
0048     auto execute_thread = BlahExecutor{params};
0049     static KernelLauncher<decltype(execute_thread)> const launch("blah");
0050     launch_kernel(count, StreamId{}, execute_thread);
0051  }
0052  * \endcode
0053  */
0054 template<class F>
0055 class KernelLauncher
0056 {
0057     static_assert(
0058         is_launchable_v<F>,
0059         R"(Launched action must be a trivially copyable function object)");
0060
0061   public:
0062     // Create a launcher from a label
0063     explicit inline KernelLauncher(std::string&& name);
0064
0065     // Launch a kernel for a thread range
0066     inline void operator()(Range<ThreadId> threads,
0067                            StreamId stream_id,
0068                            F const& execute_thread) const;
0069
0070     // Launch a kernel with a custom number of threads
0071     inline void operator()(size_type num_threads,
0072                            StreamId stream_id,
0073                            F const& execute_thread) const;
0074
0075     //! Name of this kernel
0076     std::string const& name() const { return name_; }
0077
0078   private:
0079     std::string name_;
0080     KernelParamCalculator calc_launch_params_;
0081 };
0082
0083 //---------------------------------------------------------------------------//
0084 // INLINE FUNCTIONS
0085 //---------------------------------------------------------------------------//
0086 /*!
0087  * Create a launcher from a label.
0088  */
0089 template<class F>
0090 KernelLauncher<F>::KernelLauncher(std::string&& name)
0091     : name_{std::move(name)}
0092     , calc_launch_params_{name_, &detail::launch_action_impl<F>}
0093 {
0094 }
0095
0096 //---------------------------------------------------------------------------//
0097 /*!
0098  * Launch a kernel for a thread range.
0099  */
0100 template<class F>
0101 void KernelLauncher<F>::operator()(
0102     Range<ThreadId> threads, StreamId stream_id, F const& execute_thread) const
0103 {
0104     if (!threads.empty())
0105     {
0106         ScopedProfiling profile_this_{name_};
0107         using StreamT = CELER_DEVICE_API_SYMBOL(Stream_t);
0108         StreamT stream = stream_id
0109                              ? celeritas::device().stream(stream_id).get()
0110                              : nullptr;
0111         auto config = calc_launch_params_(threads.size());
0112         detail::launch_action_impl<F>
0113             <<<config.blocks_per_grid, config.threads_per_block, 0, stream>>>(
0114                 threads, execute_thread);
0115     }
0116 }
0117
0118 //---------------------------------------------------------------------------//
0119 /*!
0120  * Launch a kernel with a custom number of threads.
0121  *
0122  * The launch arguments have the same ordering as CUDA/HIP kernel launch
0123  * arguments.
0124  *
0125  * \param num_threads Total number of active consecutive threads
0126  * \param stream_id Execute the kernel on this device stream
0127  * \param execute_thread Call the given functor with the thread ID
0128  */
0129 template<class F>
0130 void KernelLauncher<F>::operator()(
0131     size_type num_threads, StreamId stream_id, F const& execute_thread) const
0132 {
0133     (*this)(range(ThreadId{num_threads}), stream_id, execute_thread);
0134 }
0135
0136 //---------------------------------------------------------------------------//
0137 }  // namespace celeritas