Arcane  v3.16.8.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
core/RunCommandLaunchInfo.h
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2025 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
4// See the top-level COPYRIGHT file for details.
5// SPDX-License-Identifier: Apache-2.0
6//-----------------------------------------------------------------------------
7/*---------------------------------------------------------------------------*/
8/* RunCommandLaunchInfo.h (C) 2000-2025 */
9/* */
10/* Informations pour l'exécution d'une 'RunCommand'. */
11/*---------------------------------------------------------------------------*/
12#ifndef ARCANE_ACCELERATOR_CORE_RUNCOMMANDLAUNCHINFO_H
13#define ARCANE_ACCELERATOR_CORE_RUNCOMMANDLAUNCHINFO_H
14/*---------------------------------------------------------------------------*/
15/*---------------------------------------------------------------------------*/
16
18
19#include "arcane/accelerator/core/KernelLaunchArgs.h"
20
21/*---------------------------------------------------------------------------*/
22/*---------------------------------------------------------------------------*/
23
24namespace Arcane::Accelerator::impl
25{
26
27/*---------------------------------------------------------------------------*/
28/*---------------------------------------------------------------------------*/
29/*!
30 * \internal
31 * \brief Object temporaire pour conserver les informations d'exécution d'une
32 * commande et regrouper les tests.
33 */
34class ARCANE_ACCELERATOR_CORE_EXPORT RunCommandLaunchInfo
35{
36 // Les fonctions suivantes permettent de lancer les kernels.
37 template <typename SyclKernel, typename Lambda, typename LambdaArgs, typename... ReducerArgs>
38 friend void _applyKernelSYCL(impl::RunCommandLaunchInfo& launch_info, SyclKernel kernel, Lambda& func,
39 const LambdaArgs& args, const ReducerArgs&... reducer_args);
40 template <typename CudaKernel, typename Lambda, typename LambdaArgs, typename... RemainingArgs>
41 friend void _applyKernelCUDA(impl::RunCommandLaunchInfo& launch_info, const CudaKernel& kernel, Lambda& func,
42 const LambdaArgs& args, [[maybe_unused]] const RemainingArgs&... other_args);
43 template <typename HipKernel, typename Lambda, typename LambdaArgs, typename... RemainingArgs>
44 friend void _applyKernelHIP(impl::RunCommandLaunchInfo& launch_info, const HipKernel& kernel, const Lambda& func,
45 const LambdaArgs& args, [[maybe_unused]] const RemainingArgs&... other_args);
46
47 public:
48
49 using ThreadBlockInfo = KernelLaunchArgs;
50
51 public:
52
53 RunCommandLaunchInfo(RunCommand& command, Int64 total_loop_size);
54 ~RunCommandLaunchInfo();
55 RunCommandLaunchInfo(const RunCommandLaunchInfo&) = delete;
56 RunCommandLaunchInfo operator=(const RunCommandLaunchInfo&) = delete;
57
58 public:
59
60 eExecutionPolicy executionPolicy() const { return m_exec_policy; }
61
62 /*!
63 * \brief Indique qu'on commence l'exécution de la commande.
64 *
65 * Doit toujours être appelé avant de lancer la commande pour être
66 * sur que cette méthode est appelée en cas d'exception.
67 */
68 void beginExecute();
69
70 /*!
71 * \brief Signale la fin de l'exécution.
72 *
73 * Si la file associée à la commande est asynchrone, la commande
74 * peut continuer à s'exécuter après cet appel.
75 */
76 void endExecute();
77
78 /*!
79 * \brief Informations sur le nombre de block/thread/grille du noyau à lancer.
80 *
81 * Cette valeur n'est valide que pour si la commande est associée à un accélérateur.
82 */
83 KernelLaunchArgs kernelLaunchArgs() const { return m_kernel_launch_args; }
84
85 //! Calcule et retourne les informations pour les boucles multi-thread
86 ParallelLoopOptions computeParallelLoopOptions() const;
87
88 /*!
89 * \brief Informations d'exécution de la boucle.
90 *
91 * Ces informations ne sont valides que si executionPolicy()==eExecutionPolicy::Thread
92 * et si beginExecute() a été appelé.
93 */
94 const ForLoopRunInfo& loopRunInfo() const { return m_loop_run_info; }
95
96 //! Taille totale de la boucle
97 Int64 totalLoopSize() const { return m_total_loop_size; }
98
99 private:
100
101 RunCommand& m_command;
102 bool m_has_exec_begun = false;
103 bool m_is_notify_end_kernel_done = false;
105 KernelLaunchArgs m_kernel_launch_args;
106 ForLoopRunInfo m_loop_run_info;
107 Int64 m_total_loop_size = 0;
108 impl::RunQueueImpl* m_queue_impl = nullptr;
109
110 private:
111
112 KernelLaunchArgs _threadBlockInfo(const void* func, Int32 shared_memory_size) const;
113 NativeStream _internalNativeStream();
114 void _doEndKernelLaunch();
115 KernelLaunchArgs _computeKernelLaunchArgs() const;
116
117 private:
118
119 void _computeLoopRunInfo();
120
121 // Pour SYCL: enregistre l'évènement associé à la dernière commande de la file
122 // \a sycl_event_ptr est de type 'sycl::event*'.
123 void _addSyclEvent(void* sycl_event_ptr);
124};
125
126/*---------------------------------------------------------------------------*/
127/*---------------------------------------------------------------------------*/
128
129} // End namespace Arcane::Accelerator::impl
130
131/*---------------------------------------------------------------------------*/
132/*---------------------------------------------------------------------------*/
133
134#endif
Classes, Types et macros pour gérer la concurrence.
Gestion d'une commande sur accélérateur.
Type opaque pour encapsuler une 'stream' native.
Int64 totalLoopSize() const
Taille totale de la boucle.
void beginExecute()
Indique qu'on commence l'exécution de la commande.
KernelLaunchArgs kernelLaunchArgs() const
Informations sur le nombre de block/thread/grille du noyau à lancer.
void endExecute()
Signale la fin de l'exécution.
friend void _applyKernelCUDA(impl::RunCommandLaunchInfo &launch_info, const CudaKernel &kernel, Lambda &func, const LambdaArgs &args, const RemainingArgs &... other_args)
Fonction générique pour exécuter un kernel CUDA.
const ForLoopRunInfo & loopRunInfo() const
Informations d'exécution de la boucle.
friend void _applyKernelHIP(impl::RunCommandLaunchInfo &launch_info, const HipKernel &kernel, const Lambda &func, const LambdaArgs &args, const RemainingArgs &... other_args)
Fonction générique pour exécuter un kernel HIP.
Informations d'exécution d'une boucle.
Options d'exécution d'une boucle parallèle en multi-thread.
eExecutionPolicy
Politique d'exécution pour un Runner.
@ Sequential
Politique d'exécution séquentielle.
std::int64_t Int64
Type entier signé sur 64 bits.
std::int32_t Int32
Type entier signé sur 32 bits.