Arcane  v3.15.3.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
GenericScanner.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/* GenericScanner.h (C) 2000-2024 */
9/* */
10/* Algorithme de 'scan' pour les accélérateurs. */
11/*---------------------------------------------------------------------------*/
12#ifndef ARCANE_ACCELERATOR_GENERICSCANNER_H
13#define ARCANE_ACCELERATOR_GENERICSCANNER_H
14/*---------------------------------------------------------------------------*/
15/*---------------------------------------------------------------------------*/
16
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
19
20#include "arcane/utils/NumArray.h"
21
22#include "arcane/accelerator/core/RunQueue.h"
23
24#include "arcane/accelerator/AcceleratorGlobal.h"
25#include "arcane/accelerator/CommonUtils.h"
26#include "arcane/accelerator/RunCommandLaunchInfo.h"
28#include "arcane/accelerator/ScanImpl.h"
29#include "arcane/accelerator/MultiThreadAlgo.h"
30
31/*---------------------------------------------------------------------------*/
32/*---------------------------------------------------------------------------*/
33
34namespace Arcane::Accelerator::impl
35{
36
37/*---------------------------------------------------------------------------*/
38/*---------------------------------------------------------------------------*/
39/*!
40 * \internal
41 * \brief Classe pour effectuer un scan exlusif ou inclusif avec un opérateur spécifique.
42 */
44{
45 public:
46
47 explicit ScannerImpl(const RunQueue& queue)
48 : m_queue(queue)
49 {}
50
51 public:
52
53 template <bool IsExclusive, typename InputIterator, typename OutputIterator,
54 typename Operator, typename DataType>
56 DataType init_value, Operator op, const TraceInfo& trace_info)
57 {
58 RunCommand command = makeCommand(m_queue);
59 command << trace_info;
61 launch_info.beginExecute();
63 switch (exec_policy) {
64#if defined(ARCANE_COMPILING_CUDA)
66 size_t temp_storage_size = 0;
67 cudaStream_t stream = impl::CudaUtils::toNativeStream(&m_queue);
68 // Premier appel pour connaitre la taille pour l'allocation
69 if constexpr (IsExclusive)
70 ARCANE_CHECK_CUDA(::cub::DeviceScan::ExclusiveScan(nullptr, temp_storage_size,
72 else
73 ARCANE_CHECK_CUDA(::cub::DeviceScan::InclusiveScan(nullptr, temp_storage_size,
74 input_data, output_data, op, nb_item, stream));
75 void* temp_storage = m_storage.allocate(temp_storage_size);
76 if constexpr (IsExclusive)
77 ARCANE_CHECK_CUDA(::cub::DeviceScan::ExclusiveScan(temp_storage, temp_storage_size,
79 else
80 ARCANE_CHECK_CUDA(::cub::DeviceScan::InclusiveScan(temp_storage, temp_storage_size,
81 input_data, output_data, op, nb_item, stream));
82 } break;
83#endif
84#if defined(ARCANE_COMPILING_HIP)
86 size_t temp_storage_size = 0;
87 // Premier appel pour connaitre la taille pour l'allocation
88 hipStream_t stream = impl::HipUtils::toNativeStream(&m_queue);
89 if constexpr (IsExclusive)
90 ARCANE_CHECK_HIP(rocprim::exclusive_scan(nullptr, temp_storage_size, input_data, output_data,
91 init_value, nb_item, op, stream));
92 else
93 ARCANE_CHECK_HIP(rocprim::inclusive_scan(nullptr, temp_storage_size, input_data, output_data,
94 nb_item, op, stream));
95 void* temp_storage = m_storage.allocate(temp_storage_size);
96 if constexpr (IsExclusive)
98 init_value, nb_item, op, stream));
99 else
101 nb_item, op, stream));
102 } break;
103#endif
104#if defined(ARCANE_COMPILING_SYCL)
106#if defined(ARCANE_USE_SCAN_ONEDPL) && defined(__INTEL_LLVM_COMPILER)
107 sycl::queue queue = impl::SyclUtils::toNativeStream(&m_queue);
108 auto policy = oneapi::dpl::execution::make_device_policy(queue);
109 if constexpr (IsExclusive) {
110 oneapi::dpl::exclusive_scan(policy, input_data, input_data + nb_item, output_data, init_value, op);
111 }
112 else {
113 oneapi::dpl::inclusive_scan(policy, input_data, input_data + nb_item, output_data, op);
114 }
115#else
121 {
122 auto command = makeCommand(m_queue);
123 command << RUNCOMMAND_LOOP1(iter, nb_item)
124 {
125 auto [i] = iter();
126 in_data[i] = input_data[i];
127 };
128 }
129 m_queue.barrier();
131 scanner.doScan(m_queue, in_data, out_data, init_value);
132 {
133 auto command = makeCommand(m_queue);
134 command << RUNCOMMAND_LOOP1(iter, nb_item)
135 {
136 auto [i] = iter();
137 output_data[i] = out_data[i];
138 };
139 }
140 m_queue.barrier();
141#endif
142 } break;
143#endif
145 // Si le nombre de valeurs est 1 on utilise la version séquentielle.
146 // TODO: il serait judicieux de faire cela aussi pour des valeurs plus importantes
147 // car en général sur les petites boucles le multi-threading est contre productif.
148 if (nb_item > 1) {
150 scanner.doScan<IsExclusive, DataType>(launch_info.loopRunInfo(), nb_item, input_data, output_data, init_value, op);
151 break;
152 }
153 [[fallthrough]];
155 DataType sum = init_value;
156 for (Int32 i = 0; i < nb_item; ++i) {
157 DataType v = *input_data;
158 if constexpr (IsExclusive) {
159 *output_data = sum;
160 sum = op(v, sum);
161 }
162 else {
163 sum = op(v, sum);
164 *output_data = sum;
165 }
166 ++input_data;
167 ++output_data;
168 }
169 } break;
170 default:
171 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
172 }
173 launch_info.endExecute();
174 }
175
176 private:
177
178 RunQueue m_queue;
179 GenericDeviceStorage m_storage;
180};
181
182/*---------------------------------------------------------------------------*/
183/*---------------------------------------------------------------------------*/
184
185} // namespace Arcane::Accelerator::impl
186
187namespace Arcane::Accelerator
188{
189
190/*---------------------------------------------------------------------------*/
191/*---------------------------------------------------------------------------*/
192/*!
193 * \brief Algorithmes de scan exclusif ou inclusif sur accélérateurs.
194 *
195 * Voir https://en.wikipedia.org/wiki/Prefix_sum.
196 *
197 * Dans les méthodes suivantes, l'argument \a queue ne doit pas être nul.
198 */
199template <typename DataType>
200class Scanner
201{
202 public:
203
204 //! Somme exclusive
205 static void exclusiveSum(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
206 {
208 }
209 //! Minimum exclusif
210 static void exclusiveMin(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
211 {
212 _applyArray<true>(queue, input, output, ScannerMinOperator<DataType>{});
213 }
214 //! Maximum exclusif
215 static void exclusiveMax(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
216 {
217 _applyArray<true>(queue, input, output, ScannerMaxOperator<DataType>{});
218 }
219 //! Somme inclusive
220 static void inclusiveSum(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
221 {
222 _applyArray<false>(queue, input, output, ScannerSumOperator<DataType>{});
223 }
224 //! Minimum inclusif
225 static void inclusiveMin(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
226 {
227 _applyArray<false>(queue, input, output, ScannerMinOperator<DataType>{});
228 }
229 //! Maximum inclusif
230 static void inclusiveMax(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
231 {
232 _applyArray<false>(queue, input, output, ScannerMaxOperator<DataType>{});
233 }
234
235 private:
236
237 template <bool IsExclusive, typename Operator>
238 static void _applyArray(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output, const Operator& op)
239 {
241 impl::ScannerImpl scanner(*queue);
242 const Int32 nb_item = input.size();
243 if (output.size() != nb_item)
244 ARCANE_FATAL("Sizes are not equals: input={0} output={1}", nb_item, output.size());
245 const DataType* input_data = input.data();
246 DataType* output_data = output.data();
247 DataType init_value = op.defaultValue();
248 scanner.apply<IsExclusive>(nb_item, input_data, output_data, init_value, op, TraceInfo{});
249 }
250};
251
252/*---------------------------------------------------------------------------*/
253/*---------------------------------------------------------------------------*/
254/*!
255 * \brief Algorithmes de scan exclusif ou inclusif sur accélérateurs.
256 *
257 * Voir https://en.wikipedia.org/wiki/Prefix_sum.
258 *
259 * Dans les méthodes de scan, les valeurs entre les entrées et les sorties
260 * ne doivent pas se chevaucher.
261 */
262class GenericScanner
263{
264 public:
265
266 /*!
267 * \brief Itérateur sur une lambda pour positionner une valeur via un index.
268 */
269 template <typename DataType, typename SetterLambda>
271 {
272 public:
273
274 //! Permet de positionner un élément de l'itérateur de sortie
275 class Setter
276 {
277 public:
278
279 ARCCORE_HOST_DEVICE explicit Setter(const SetterLambda& s, Int32 index)
280 : m_index(index)
281 , m_lambda(s)
282 {}
283 ARCCORE_HOST_DEVICE void operator=(const DataType& value)
284 {
285 m_lambda(m_index, value);
286 }
287
288 public:
289
290 Int32 m_index = 0;
291 SetterLambda m_lambda;
292 };
293
294 using value_type = DataType;
295 using iterator_category = std::random_access_iterator_tag;
296 using reference = Setter;
298 using pointer = void;
300
301 public:
302
304 : m_lambda(s)
305 {}
306 ARCCORE_HOST_DEVICE explicit SetterLambdaIterator(const SetterLambda& s, Int32 v)
307 : m_index(v)
308 , m_lambda(s)
309 {}
310
311 public:
312
313 ARCCORE_HOST_DEVICE ThatClass& operator++()
314 {
315 ++m_index;
316 return (*this);
317 }
318 ARCCORE_HOST_DEVICE friend ThatClass operator+(const ThatClass& iter, Int32 x)
319 {
320 return ThatClass(iter.m_lambda, iter.m_index + x);
321 }
322 ARCCORE_HOST_DEVICE friend ThatClass operator+(Int32 x, const ThatClass& iter)
323 {
324 return ThatClass(iter.m_lambda, iter.m_index + x);
325 }
326 ARCCORE_HOST_DEVICE friend bool operator<(const ThatClass& iter1, const ThatClass& iter2)
327 {
328 return iter1.m_index < iter2.m_index;
329 }
330 ARCCORE_HOST_DEVICE ThatClass operator-(Int32 x)
331 {
332 return ThatClass(m_lambda, m_index - x);
333 }
334 ARCCORE_HOST_DEVICE Int32 operator-(const ThatClass& x) const
335 {
336 return m_index - x.m_index;
337 }
338 ARCCORE_HOST_DEVICE reference operator*() const
339 {
340 return Setter(m_lambda, m_index);
341 }
342 ARCCORE_HOST_DEVICE reference operator[](Int32 x) const { return Setter(m_lambda, m_index + x); }
343 ARCCORE_HOST_DEVICE friend bool operator!=(const ThatClass& a, const ThatClass& b)
344 {
345 return a.m_index != b.m_index;
346 }
347
348 private:
349
350 Int32 m_index = 0;
351 SetterLambda m_lambda;
352 };
353
354 public:
355
356 explicit GenericScanner(const RunQueue& queue)
357 : m_queue(queue)
358 {}
359
360 public:
361
362 template <typename DataType, typename GetterLambda, typename SetterLambda, typename Operator>
363 void applyWithIndexExclusive(Int32 nb_value, const DataType& initial_value,
364 const GetterLambda& getter_lambda,
365 const SetterLambda& setter_lambda,
366 const Operator& op_lambda,
367 const TraceInfo& trace_info = TraceInfo())
368 {
369 _applyWithIndex<true>(nb_value, initial_value, getter_lambda, setter_lambda, op_lambda, trace_info);
370 }
371
372 template <typename DataType, typename GetterLambda, typename SetterLambda, typename Operator>
373 void applyWithIndexInclusive(Int32 nb_value, const DataType& initial_value,
374 const GetterLambda& getter_lambda,
375 const SetterLambda& setter_lambda,
376 const Operator& op_lambda,
377 const TraceInfo& trace_info = TraceInfo())
378 {
379 _applyWithIndex<false>(nb_value, initial_value, getter_lambda, setter_lambda, op_lambda, trace_info);
380 }
381
382 template <typename InputDataType, typename OutputDataType, typename Operator>
383 void applyExclusive(const OutputDataType& initial_value,
384 SmallSpan<const InputDataType> input,
385 SmallSpan<OutputDataType> output,
386 const Operator& op_lambda,
387 const TraceInfo& trace_info = TraceInfo())
388 {
389 _apply<true>(initial_value, input, output, op_lambda, trace_info);
390 }
391
392 template <typename InputDataType, typename OutputDataType, typename Operator>
393 void applyInclusive(const OutputDataType& initial_value,
394 SmallSpan<const InputDataType> input,
395 SmallSpan<OutputDataType> output,
396 const Operator& op_lambda,
397 const TraceInfo& trace_info = TraceInfo())
398 {
399 _apply<false>(initial_value, input, output, op_lambda, trace_info);
400 }
401
402 private:
403
404 template <bool IsExclusive, typename DataType, typename GetterLambda, typename SetterLambda, typename Operator>
405 void _applyWithIndex(Int32 nb_value, const DataType& initial_value,
406 const GetterLambda& getter_lambda,
407 const SetterLambda& setter_lambda,
408 const Operator& op_lambda,
409 const TraceInfo& trace_info)
410 {
411 impl::GetterLambdaIterator<DataType, GetterLambda> input_iter(getter_lambda);
412 SetterLambdaIterator<DataType, SetterLambda> output_iter(setter_lambda);
413 impl::ScannerImpl scanner(m_queue);
414 scanner.apply<IsExclusive>(nb_value, input_iter, output_iter, initial_value, op_lambda, trace_info);
415 }
416
417 template <bool IsExclusive, typename InputDataType, typename OutputDataType, typename Operator>
418 void _apply(const OutputDataType& initial_value,
419 SmallSpan<const InputDataType> input,
420 SmallSpan<OutputDataType> output,
421 const Operator& op,
422 const TraceInfo& trace_info = TraceInfo())
423 {
424 const Int32 nb_item = input.size();
425 if (output.size() != nb_item)
426 ARCANE_FATAL("Sizes are not equals: input={0} output={1}", nb_item, output.size());
427 auto* input_data = input.data();
428 auto* output_data = output.data();
429 impl::ScannerImpl scanner(m_queue);
430 scanner.apply<IsExclusive>(nb_item, input_data, output_data, initial_value, op, trace_info);
431 }
432
433 private:
434
435 RunQueue m_queue;
436};
437
438/*---------------------------------------------------------------------------*/
439/*---------------------------------------------------------------------------*/
440
441} // namespace Arcane::Accelerator
442
443/*---------------------------------------------------------------------------*/
444/*---------------------------------------------------------------------------*/
445
446#endif
447
448/*---------------------------------------------------------------------------*/
449/*---------------------------------------------------------------------------*/
#define ARCANE_CHECK_POINTER(ptr)
Macro retournant le pointeur ptr s'il est non nul ou lancant une exception s'il est nul.
#define ARCANE_FATAL(...)
Macro envoyant une exception FatalErrorException.
Types et macros pour gérer les boucles sur les accélérateurs.
#define RUNCOMMAND_LOOP1(iter_name, x1,...)
Boucle sur accélérateur avec arguments supplémentaires pour les réductions.
Permet de positionner un élément de l'itérateur de sortie.
Itérateur sur une lambda pour positionner une valeur via un index.
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
Algorithmes avancée en mode multi-thread.
Référence à une instance.
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.