Arcane  v3.15.0.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
GenericReducer.h
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2024 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/* GenericReducer.h (C) 2000-2024 */
9/* */
10/* Gestion des réductions pour les accélérateurs. */
11/*---------------------------------------------------------------------------*/
12#ifndef ARCANE_ACCELERATOR_GENERICREDUCER_H
13#define ARCANE_ACCELERATOR_GENERICREDUCER_H
14/*---------------------------------------------------------------------------*/
15/*---------------------------------------------------------------------------*/
16
17#include "arcane/utils/NumArray.h"
18#include "arcane/utils/FatalErrorException.h"
19
20#include "arcane/accelerator/core/RunQueue.h"
22
23/*---------------------------------------------------------------------------*/
24/*---------------------------------------------------------------------------*/
25
26namespace Arcane::Accelerator::impl
27{
28template <typename DataType>
29class GenericReducerIf;
30
31/*---------------------------------------------------------------------------*/
32/*---------------------------------------------------------------------------*/
33// Classe pour déterminer l'instance de 'Reducer2' à utiliser en fonction de l'opérateur.
34// A spécialiser.
35template <typename DataType, typename Operator>
37
38template <typename DataType>
40{
41 public:
42
43 using ReducerType = ReducerMax2<DataType>;
44};
45template <typename DataType>
47{
48 public:
49
50 using ReducerType = ReducerMin2<DataType>;
51};
52template <typename DataType>
54{
55 public:
56
57 using ReducerType = ReducerSum2<DataType>;
58};
59
60/*---------------------------------------------------------------------------*/
61/*---------------------------------------------------------------------------*/
62/*!
63 * \internal
64 * \brief Classe de base pour effectuer une réduction.
65 *
66 * Contient les arguments nécessaires pour effectuer une réduction.
67 */
68template <typename DataType>
70{
71 friend class GenericReducerIf<DataType>;
72
73 public:
74
75 GenericReducerBase(const RunQueue& queue)
76 : m_queue(queue)
77 {}
78
79 protected:
80
81 DataType _reducedValue() const
82 {
83 m_queue.barrier();
84 return m_host_reduce_storage[0];
85 }
86
87 void _allocate()
88 {
89 eMemoryRessource r = eMemoryRessource::HostPinned;
90 if (m_host_reduce_storage.memoryRessource() != r)
91 m_host_reduce_storage = NumArray<DataType, MDDim1>(r);
92 m_host_reduce_storage.resize(1);
93 }
94
95 protected:
96
97 RunQueue m_queue;
98 GenericDeviceStorage m_algo_storage;
99 DeviceStorage<DataType> m_device_reduce_storage;
100 NumArray<DataType, MDDim1> m_host_reduce_storage;
101};
102
103/*---------------------------------------------------------------------------*/
104/*---------------------------------------------------------------------------*/
105/*!
106 * \internal
107 * \brief Classe pour effectuer un partitionnement d'une liste.
108 *
109 * La liste est partitionnée en deux listes.
110 *
111 * \a DataType est le type de donnée.
112 */
113template <typename DataType>
115{
116 // TODO: Faire le malloc sur le device associé à la queue.
117 // et aussi regarder si on peut utiliser mallocAsync().
118
119 public:
120
121 template <typename InputIterator, typename ReduceOperator>
122 void apply(GenericReducerBase<DataType>& s, Int32 nb_item, const DataType& init_value,
123 InputIterator input_iter, ReduceOperator reduce_op, const TraceInfo& trace_info)
124 {
125 RunQueue& queue = s.m_queue;
126 RunCommand command = makeCommand(queue);
127 command << trace_info;
128 impl::RunCommandLaunchInfo launch_info(command, nb_item);
129 launch_info.beginExecute();
130 eExecutionPolicy exec_policy = queue.executionPolicy();
131 switch (exec_policy) {
132#if defined(ARCANE_COMPILING_CUDA)
134 size_t temp_storage_size = 0;
135 cudaStream_t stream = impl::CudaUtils::toNativeStream(queue);
136 DataType* reduced_value_ptr = nullptr;
137 // Premier appel pour connaitre la taille pour l'allocation
138 ARCANE_CHECK_CUDA(::cub::DeviceReduce::Reduce(nullptr, temp_storage_size, input_iter, reduced_value_ptr,
139 nb_item, reduce_op, init_value, stream));
140
141 s.m_algo_storage.allocate(temp_storage_size);
142 reduced_value_ptr = s.m_device_reduce_storage.allocate();
143 ARCANE_CHECK_CUDA(::cub::DeviceReduce::Reduce(s.m_algo_storage.address(), temp_storage_size,
144 input_iter, reduced_value_ptr, nb_item,
145 reduce_op, init_value, stream));
146 s.m_device_reduce_storage.copyToAsync(s.m_host_reduce_storage, queue);
147 } break;
148#endif
149#if defined(ARCANE_COMPILING_HIP)
151 size_t temp_storage_size = 0;
152 hipStream_t stream = impl::HipUtils::toNativeStream(queue);
153 DataType* reduced_value_ptr = nullptr;
154 // Premier appel pour connaitre la taille pour l'allocation
155 ARCANE_CHECK_HIP(rocprim::reduce(nullptr, temp_storage_size, input_iter, reduced_value_ptr, init_value,
156 nb_item, reduce_op, stream));
157
158 s.m_algo_storage.allocate(temp_storage_size);
159 reduced_value_ptr = s.m_device_reduce_storage.allocate();
160
161 ARCANE_CHECK_HIP(rocprim::reduce(s.m_algo_storage.address(), temp_storage_size, input_iter, reduced_value_ptr, init_value,
162 nb_item, reduce_op, stream));
163 s.m_device_reduce_storage.copyToAsync(s.m_host_reduce_storage, queue);
164 } break;
165#endif
166#if defined(ARCANE_COMPILING_SYCL)
168 {
169 RunCommand command2 = makeCommand(queue);
171 ReducerType reducer(command2);
172 command2 << RUNCOMMAND_LOOP1(iter, nb_item, reducer)
173 {
174 auto [i] = iter();
175 reducer.combine(input_iter[i]);
176 };
177 queue.barrier();
178 s.m_host_reduce_storage[0] = reducer.reducedValue();
179 }
180 } break;
181#endif
183 // Pas encore implémenté en multi-thread
184 [[fallthrough]];
186 DataType reduced_value = init_value;
187 for (Int32 i = 0; i < nb_item; ++i) {
188 reduced_value = reduce_op(reduced_value, *input_iter);
189 ++input_iter;
190 }
191 s.m_host_reduce_storage[0] = reduced_value;
192 } break;
193 default:
194 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
195 }
196 launch_info.endExecute();
197 }
198};
199
200/*---------------------------------------------------------------------------*/
201/*---------------------------------------------------------------------------*/
202
203} // namespace Arcane::Accelerator::impl
204
205/*---------------------------------------------------------------------------*/
206/*---------------------------------------------------------------------------*/
207
208namespace Arcane::Accelerator
209{
210
211/*---------------------------------------------------------------------------*/
212/*---------------------------------------------------------------------------*/
213/*!
214 * \brief Algorithme générique de réduction sur accélérateur.
215 *
216 * La réduction se fait via les appels à applyMin(), applyMax(), applySum(),
217 * applyMinWithIndex(), applyMaxWithIndex() ou applySumWithIndex(). Ces
218 * méthodes sont asynchrones. Après réduction, il est possible récupérer la
219 * valeur réduite via reducedValue(). L'appel à reducedValue() bloque tant
220 * que la réduction n'est pas terminée.
221 *
222 * Les instances de cette classe peuvent être utilisées plusieurs fois.
223 *
224 * Voici un exemple pour calculer la somme d'un tableau de 50 éléments:
225 *
226 * \code
227 * using namespace Arcane;
228 * const Int32 nb_value(50);
229 * Arcane::NumArray<Real, MDDim1> t1(nv_value);
230 * Arcane::SmallSpan<const Real> t1_view(t1);
231 * Arcane::RunQueue queue = ...;
232 * Arcane::Accelerator::GenericReducer<Real> reducer(queue);
233 *
234 * // Calcul direct
235 * reducer.applySum(t1_view);
236 * std::cout << "Sum is '" << reducer.reducedValue() << "\n";
237 *
238 * // Calcul avec lambda
239 * auto getter_func = [=] ARCCORE_HOST_DEVICE(Int32 index) -> Real
240 * {
241 * return t1_view[index];
242 * }
243 * reducer.applySumWithIndex(nb_value,getter_func);
244 * std::cout << "Sum is '" << reducer.reducedValue() << "\n";
245 * \endcode
246 */
247template <typename DataType>
249: private impl::GenericReducerBase<DataType>
250{
251 public:
252
253 explicit GenericReducer(const RunQueue& queue)
255 {
256 this->_allocate();
257 }
258
259 public:
260
261 //! Applique une réduction 'Min' sur les valeurs \a values
262 void applyMin(SmallSpan<const DataType> values, const TraceInfo& trace_info = TraceInfo())
263 {
264 _apply(values.size(), values.data(), impl::MinOperator<DataType>{}, trace_info);
265 }
266
267 //! Applique une réduction 'Max' sur les valeurs \a values
268 void applyMax(SmallSpan<const DataType> values, const TraceInfo& trace_info = TraceInfo())
269 {
270 _apply(values.size(), values.data(), impl::MaxOperator<DataType>{}, trace_info);
271 }
272
273 //! Applique une réduction 'Somme' sur les valeurs \a values
274 void applySum(SmallSpan<const DataType> values, const TraceInfo& trace_info = TraceInfo())
275 {
276 _apply(values.size(), values.data(), impl::SumOperator<DataType>{}, trace_info);
277 }
278
279 //! Applique une réduction 'Min' sur les valeurs sélectionnées par \a select_lambda
280 template <typename SelectLambda>
281 void applyMinWithIndex(Int32 nb_value, const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
282 {
283 _applyWithIndex(nb_value, select_lambda, impl::MinOperator<DataType>{}, trace_info);
284 }
285
286 //! Applique une réduction 'Max' sur les valeurs sélectionnées par \a select_lambda
287 template <typename SelectLambda>
288 void applyMaxWithIndex(Int32 nb_value, const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
289 {
290 _applyWithIndex(nb_value, select_lambda, impl::MaxOperator<DataType>{}, trace_info);
291 }
292
293 //! Applique une réduction 'Somme' sur les valeurs sélectionnées par \a select_lambda
294 template <typename SelectLambda>
295 void applySumWithIndex(Int32 nb_value, const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
296 {
297 _applyWithIndex(nb_value, select_lambda, impl::SumOperator<DataType>{}, trace_info);
298 }
299
300 //! Valeur de la réduction
301 DataType reducedValue()
302 {
303 m_is_already_called = false;
304 return this->_reducedValue();
305 }
306
307 private:
308
309 bool m_is_already_called = false;
310
311 private:
312
313 template <typename InputIterator, typename ReduceOperator>
314 void _apply(Int32 nb_value, InputIterator input_iter, ReduceOperator reduce_op, const TraceInfo& trace_info)
315 {
316 _setCalled();
317 impl::GenericReducerBase<DataType>* base_ptr = this;
319 DataType init_value = reduce_op.defaultValue();
320 gf.apply(*base_ptr, nb_value, init_value, input_iter, reduce_op, trace_info);
321 }
322
323 template <typename GetterLambda, typename ReduceOperator>
324 void _applyWithIndex(Int32 nb_value, const GetterLambda& getter_lambda,
325 ReduceOperator reduce_op, const TraceInfo& trace_info)
326 {
327 _setCalled();
328 impl::GenericReducerBase<DataType>* base_ptr = this;
329 impl::GenericReducerIf<DataType> gf;
330 impl::GetterLambdaIterator<DataType, GetterLambda> input_iter(getter_lambda);
331 DataType init_value = reduce_op.defaultValue();
332 gf.apply(*base_ptr, nb_value, init_value, input_iter, reduce_op, trace_info);
333 }
334
335 void _setCalled()
336 {
337 if (m_is_already_called)
338 ARCANE_FATAL("apply() has already been called for this instance");
339 m_is_already_called = true;
340 }
341};
342
343/*---------------------------------------------------------------------------*/
344/*---------------------------------------------------------------------------*/
345
346} // namespace Arcane::Accelerator
347
348/*---------------------------------------------------------------------------*/
349/*---------------------------------------------------------------------------*/
350
351#endif
352
353/*---------------------------------------------------------------------------*/
354/*---------------------------------------------------------------------------*/
#define ARCANE_FATAL(...)
Macro envoyant une exception FatalErrorException.
Types et fonctions pour gérer les synchronisations sur les accélérateurs.
#define RUNCOMMAND_LOOP1(iter_name, x1,...)
Boucle sur accélérateur avec arguments supplémentaires pour les réductions.
Algorithme générique de réduction sur accélérateur.
DataType reducedValue()
Valeur de la réduction.
void applyMinWithIndex(Int32 nb_value, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Min' sur les valeurs sélectionnées par select_lambda.
void applyMaxWithIndex(Int32 nb_value, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Max' sur les valeurs sélectionnées par select_lambda.
void applyMax(SmallSpan< const DataType > values, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Max' sur les valeurs values.
void applyMin(SmallSpan< const DataType > values, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Min' sur les valeurs values.
void applySum(SmallSpan< const DataType > values, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Somme' sur les valeurs values.
void applySumWithIndex(Int32 nb_value, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique une réduction 'Somme' sur les valeurs sélectionnées par select_lambda.
Gestion d'une commande sur accélérateur.
File d'exécution pour un accélérateur.
void barrier() const
Bloque tant que toutes les commandes associées à la file ne sont pas terminées.
Definition RunQueue.cc:159
eExecutionPolicy executionPolicy() const
Politique d'exécution de la file.
Definition RunQueue.cc:169
Opérateur de Scan/Reduce pour le maximum.
Opérateur de Scan/Reduce pour le minimum.
void beginExecute()
Indique qu'on commence l'exécution de la commande.
void endExecute()
Signale la fin de l'exécution.
Opérateur de Scan/Reduce pour les sommes.
Tableaux multi-dimensionnels pour les types numériques accessibles sur accélérateurs.
Vue d'un tableau d'éléments de type T.
Definition Span.h:670
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:209
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:419
Espace de nom pour l'utilisation des accélérateurs.
RunCommand makeCommand(const RunQueue &run_queue)
Créé une commande associée à la file run_queue.
eExecutionPolicy
Politique d'exécution pour un Runner.
@ SYCL
Politique d'exécution utilisant l'environnement SYCL.
@ HIP
Politique d'exécution utilisant l'environnement HIP.
@ CUDA
Politique d'exécution utilisant l'environnement CUDA.
@ Sequential
Politique d'exécution séquentielle.
@ Thread
Politique d'exécution multi-thread.
eMemoryResource
Liste des ressources mémoire disponibles.
std::int32_t Int32
Type entier signé sur 32 bits.