Arcane  v3.16.0.0
Documentation développeur
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>
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/*---------------------------------------------------------------------------*/
68template <typename DataType>
69class GenericReducerBase
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/*---------------------------------------------------------------------------*/
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/*---------------------------------------------------------------------------*/
247template <typename DataType>
248class GenericReducer
249: private impl::GenericReducerBase<DataType>
250{
251 public:
252
253 explicit GenericReducer(const RunQueue& queue)
255 {
256 this->_allocate();
257 }
258
259 public:
260
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
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
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
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
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
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
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.
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.
Classe pour effectuer une réduction 'max'.
Definition Reduce.h:775
Classe pour effectuer une réduction 'min'.
Definition Reduce.h:801
Classe pour effectuer une réduction 'somme'.
Definition Reduce.h:750
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
Gère l'allocation interne sur le device pour un type donné.
Gère l'allocation interne sur le device.
Definition CommonUtils.h:95
Classe de base pour effectuer une réduction.
Classe pour effectuer un partitionnement d'une liste.
Opérateur de Scan/Reduce pour le maximum.
Opérateur de Scan/Reduce pour le minimum.
Object temporaire pour conserver les informations d'exécution d'une commande et regrouper les tests.
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:673
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:212
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:422
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.
Arcane::eMemoryResource eMemoryRessource
Typedef pour la version Arcane historique (avec 2's')
std::int32_t Int32
Type entier signé sur 32 bits.