Arcane  v3.15.3.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
GenericPartitioner.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/* GenericPartitioner.h (C) 2000-2025 */
9/* */
10/* Algorithme de partitionnement de liste. */
11/*---------------------------------------------------------------------------*/
12#ifndef ARCANE_ACCELERATOR_GENERICPARTITIONER_H
13#define ARCANE_ACCELERATOR_GENERICPARTITIONER_H
14/*---------------------------------------------------------------------------*/
15/*---------------------------------------------------------------------------*/
16
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
19#include "arcane/utils/NumArray.h"
20
21#include "arcane/accelerator/AcceleratorGlobal.h"
22#include "arcane/accelerator/core/RunQueue.h"
23#include "arcane/accelerator/CommonUtils.h"
24#include "arcane/accelerator/RunCommandLaunchInfo.h"
25#if defined(ARCANE_COMPILING_SYCL)
27#endif
28
29/*---------------------------------------------------------------------------*/
30/*---------------------------------------------------------------------------*/
31
32namespace Arcane::Accelerator::impl
33{
34
35/*---------------------------------------------------------------------------*/
36/*---------------------------------------------------------------------------*/
37/*!
38 * \internal
39 * \brief Classe de base pour effectuer un filtrage.
40 *
41 * Contient les arguments nécessaires pour effectuer le filtrage.
42 */
43class ARCANE_ACCELERATOR_EXPORT GenericPartitionerBase
44{
45 friend class GenericPartitionerIf;
46
47 public:
48
49 explicit GenericPartitionerBase(const RunQueue& queue);
50
51 protected:
52
53 Int32 _nbFirstPart() const;
54 SmallSpan<const Int32> _nbParts() const;
55 void _allocate();
56 void _resetNbPart();
57
58 protected:
59
60 RunQueue m_queue;
61 GenericDeviceStorage m_algo_storage;
62 DeviceStorage<int, 2> m_device_nb_list1_storage;
63 NumArray<Int32, MDDim1> m_host_nb_list1_storage;
64};
65
66/*---------------------------------------------------------------------------*/
67/*---------------------------------------------------------------------------*/
68/*!
69 * \internal
70 * \brief Classe pour effectuer un partitionnement d'une liste.
71 */
73{
74 public:
75
76 /*!
77 * \brief Effectue le partitionnement d'une liste en deux parties.
78 */
79 template <typename SelectLambda, typename InputIterator, typename OutputIterator>
82 {
83 RunQueue queue = s.m_queue;
85 RunCommand command = makeCommand(queue);
86 command << trace_info;
88 launch_info.beginExecute();
89 switch (exec_policy) {
90#if defined(ARCANE_COMPILING_CUDA)
92 size_t temp_storage_size = 0;
93 cudaStream_t stream = impl::CudaUtils::toNativeStream(&queue);
94 // Premier appel pour connaitre la taille pour l'allocation
95 int* nb_list1_ptr = nullptr;
96 ARCANE_CHECK_CUDA(::cub::DevicePartition::If(nullptr, temp_storage_size,
98 select_lambda, stream));
99
100 s.m_algo_storage.allocate(temp_storage_size);
101 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
102 ARCANE_CHECK_CUDA(::cub::DevicePartition::If(s.m_algo_storage.address(), temp_storage_size,
104 select_lambda, stream));
105 s.m_device_nb_list1_storage.copyToAsync(s.m_host_nb_list1_storage, queue);
106 } break;
107#endif
108#if defined(ARCANE_COMPILING_HIP)
110 size_t temp_storage_size = 0;
111 // Premier appel pour connaitre la taille pour l'allocation
112 hipStream_t stream = impl::HipUtils::toNativeStream(&queue);
113 int* nb_list1_ptr = nullptr;
114 ARCANE_CHECK_HIP(rocprim::partition(nullptr, temp_storage_size, input_iter, output_iter,
116
117 s.m_algo_storage.allocate(temp_storage_size);
118 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
119
120 ARCANE_CHECK_HIP(rocprim::partition(s.m_algo_storage.address(), temp_storage_size, input_iter, output_iter,
122 s.m_device_nb_list1_storage.copyToAsync(s.m_host_nb_list1_storage, queue);
123 } break;
124#endif
125#if defined(ARCANE_COMPILING_SYCL) && defined(__INTEL_LLVM_COMPILER)
127 // Seulement implémenté pour DPC++.
128 // Actuellement (dpc++ 2025.0), il n'y a pas l'équivalent avec SYCL de
129 // la méthode de partition de cub ou rocprim.
130 // Utilise la fonction 'stable_partition'. Cependant, cette fonction
131 // ne supporte pas si la mémoire uniquement sur accélérateur. De plus,
132 // il faut que InputIterator et OutputIterator remplissent le concept
133 // std::random_access_iterator ce qui n'est pas (notamment car il n'y a
134 // pas les opérateurs de copie ni de constructeur vide à cause de la lambda).
135 // Pour éviter tous ces problèmes, on alloue donc des tableaux temporaires
136 // pour l'appel à 'stable_sort' et on recopie les valerus en sortie.
137 using InputDataType = typename InputIterator::value_type;
138 using DataType = typename OutputIterator::value_type;
141 auto tmp_output = tmp_output_numarray.to1DSmallSpan();
142 auto tmp_select = tmp_select_numarray.to1DSmallSpan();
143 {
144 auto command = makeCommand(queue);
145 command << RUNCOMMAND_LOOP1(iter, nb_item)
146 {
147 auto [i] = iter();
148 tmp_output[i] = input_iter[i];
149 };
150 }
151 auto tmp_select_lambda = [=](Int32 i) { return tmp_select[i]; };
152 sycl::queue sycl_queue = impl::SyclUtils::toNativeStream(queue);
153 auto policy = oneapi::dpl::execution::make_device_policy(sycl_queue);
154 auto output_after = oneapi::dpl::stable_partition(policy, tmp_output.begin(), tmp_output.end(), select_lambda);
155 queue.barrier();
156 Int32 nb_list1 = (output_after - tmp_output.begin());
157 Int32 nb_list2 = nb_item - nb_list1;
158 s.m_host_nb_list1_storage[0] = nb_list1;
159 //std::cerr << "NbList1=" << nb_list1 << " NbList2=" << nb_list2 << "\n";
160 {
161 // Recopie dans output les valeurs filtrées.
162 // Pour être cohérent avec 'cub' et 'rocprim', il faut inverser l'ordre des
163 // des valeurs de la liste pour les éléments ne remplissant pas la condition.
164 // Pour cela, on fait une boucle de taille (nb_list1 + nb_list/2) et chaque
165 // itération pour i>=nb_list1 gère deux éléments.
166 auto command = makeCommand(queue);
167 Int32 nb_iter2 = (nb_list2 / 2) + (nb_list2 % 2);
168 //std::cout << "NB_ITER2=" << nb_iter2 << "\n";
169 command << RUNCOMMAND_LOOP1(iter, (nb_list1 + nb_iter2))
170 {
171 auto [i] = iter();
172 if (i >= nb_list1) {
173 // Partie de la liste pour les valeurs ne remplissant par le critère.
174 Int32 j = i - nb_list1;
175 Int32 reverse_i = (nb_item - (j + 1));
176 auto x1 = tmp_output[i];
177 auto x2 = tmp_output[reverse_i];
180 }
181 else
182 output_iter[i] = tmp_output[i];
183 };
184 }
185 queue.barrier();
186 } break;
187#endif
189 // Pas encore implémenté en multi-thread
190 [[fallthrough]];
194 for (Int32 i = 0; i < nb_item; ++i) {
195 auto v = *input_iter;
196 if (select_lambda(v)) {
197 *output_iter = v;
198 ++output_iter;
199 }
200 else {
201 --output2_iter;
202 *output2_iter = v;
203 }
204 ++input_iter;
205 }
206 Int32 nb_list1 = static_cast<Int32>(output_iter - saved_output_iter);
207 s.m_host_nb_list1_storage[0] = nb_list1;
208 } break;
209 default:
210 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
211 }
212 launch_info.endExecute();
213 }
214
215 /*!
216 * \brief Effectue le partitionnement d'une liste en trois parties.
217 */
218 template <typename Select1Lambda, typename Select2Lambda,
219 typename InputIterator, typename FirstOutputIterator,
220 typename SecondOutputIterator, typename UnselectedIterator>
228 const TraceInfo& trace_info = TraceInfo())
229 {
230 RunQueue queue = s.m_queue;
232 RunCommand command = makeCommand(queue);
233 command << trace_info;
235 launch_info.beginExecute();
236 switch (exec_policy) {
237#if defined(ARCANE_COMPILING_CUDA)
239 size_t temp_storage_size = 0;
240 cudaStream_t stream = impl::CudaUtils::toNativeStream(&queue);
241 // Premier appel pour connaitre la taille pour l'allocation
242 int* nb_list1_ptr = nullptr;
243 ARCANE_CHECK_CUDA(::cub::DevicePartition::If(nullptr, temp_storage_size,
247
248 s.m_algo_storage.allocate(temp_storage_size);
249 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
250 ARCANE_CHECK_CUDA(::cub::DevicePartition::If(s.m_algo_storage.address(), temp_storage_size,
254 s.m_device_nb_list1_storage.copyToAsync(s.m_host_nb_list1_storage, queue);
255 } break;
256#endif
257#if defined(ARCANE_COMPILING_HIP)
259 size_t temp_storage_size = 0;
260 // Premier appel pour connaitre la taille pour l'allocation
261 hipStream_t stream = impl::HipUtils::toNativeStream(&queue);
262 int* nb_list1_ptr = nullptr;
263 using namespace rocprim;
264 ARCANE_CHECK_HIP(::rocprim::partition_three_way(nullptr, temp_storage_size, input_iter, first_output_iter,
267
268 s.m_algo_storage.allocate(temp_storage_size);
269 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
270
274 s.m_device_nb_list1_storage.copyToAsync(s.m_host_nb_list1_storage, queue);
275 } break;
276#endif
278 // Pas encore implémenté en multi-thread
279 [[fallthrough]];
281 Int32 nb_first = 0;
282 Int32 nb_second = 0;
283 for (Int32 i = 0; i < nb_item; ++i) {
284 auto v = *input_iter;
285 bool is_1 = select1_lambda(v);
286 bool is_2 = select2_lambda(v);
287 if (is_1) {
290 ++nb_first;
291 }
292 else {
293 if (is_2) {
296 ++nb_second;
297 }
298 else {
299 *unselected_iter = v;
301 }
302 }
303 // Incrémenter l'itérateur à la fin car il est utilisé pour le positionnement
304 ++input_iter;
305 }
306 s.m_host_nb_list1_storage[0] = nb_first;
307 s.m_host_nb_list1_storage[1] = nb_second;
308 } break;
309 default:
310 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
311 }
312 launch_info.endExecute();
313 }
314};
315
316/*---------------------------------------------------------------------------*/
317/*---------------------------------------------------------------------------*/
318
319} // namespace Arcane::Accelerator::impl
320
321namespace Arcane::Accelerator
322{
323
324/*---------------------------------------------------------------------------*/
325/*---------------------------------------------------------------------------*/
326/*!
327 * \brief Algorithme générique de partitionnement d'une liste.
328 *
329 * Cette classe fournit des algorithmes pour partitionner une liste en deux
330 * ou trois parties selon un critère fixé par l'utilisateur.
331 */
334{
335 public:
336
337 explicit GenericPartitioner(const RunQueue& queue)
339 {
340 _allocate();
341 }
342
343 public:
344
345 /*!
346 * \brief Effectue un partitionnement d'une liste en deux parties.
347 *
348 * Le nombre de valeurs de la liste est donné par \a nb_value.
349 * Les deux fonctions lambda \a select_lambda et \a setter_lambda permettent
350 * de partitionner et de positionner les valeurs de la liste.
351 *
352 * Après exécution, il est possible de récupérer le nombre d'éléments
353 * de la première partie de la liste via la méthode \a nbFirstPart().
354 *
355 * \snippet AcceleratorPartitionerUnitTest.cc SampleListPartitionerTwoWayIndex
356 */
357 template <typename SelectLambda, typename SetterLambda>
371
372 /*!
373 * \brief Effectue un partitionnement d'une liste en deux parties.
374 *
375 * Le nombre de valeurs de la liste est donné par \a nb_value.
376 * Les valeurs en entrée sont fournies par l'itérateur \a input_iter et
377 * les valeurs en sorties par l'itérateur \a output_iterator. La fonction
378 * lambda \a select_lambda permet de sélectionner la partition utilisée :
379 * si le retour est \a true, la valeur sera dans la première partie de la liste,
380 * sinon elle sera dans la seconde. En sortie les valeurs de la deuxième
381 * partie sont rangées en ordre inverse de la liste d'entrée.
382 *
383 * Après exécution, il est possible de récupérer le nombre d'éléments
384 * de la première partie de la liste via la méthode \a nbFirstPart().
385 *
386 * \snippet AcceleratorPartitionerUnitTest.cc SampleListPartitionerTwoWayIf
387 */
388 template <typename InputIterator, typename OutputIterator, typename SelectLambda>
399
400 /*!
401 * \brief Effectue un partitionnement d'une liste en trois parties.
402 *
403 * Le nombre de valeurs de la liste est donné par \a nb_value.
404 * Les deux fonctions lambda \a select1_lambda et \a select2_lambda permettent
405 * de partitionner la liste avec l'algorithme suivant:
406 * - si select1_lambda() est vrai, la valeur sera positionnée via \a setter1_lambda,
407 * - sinon si select2_lambda() est vrai, la valeur sera positionnée via \a setter2_lambda,
408 * - sinon la valeur sera positionnée via \a unselected_setter_lambda.
409 *
410 * Les listes en sortie sont dans le même ordre qu'en entrée.
411 *
412 * Après exécution, il est possible de récupérer le nombre d'éléments
413 * de la première et de la deuxième liste la méthode \a nbParts().
414 *
415 * \snippet AcceleratorPartitionerUnitTest.cc SampleListPartitionerThreeWayIndex
416 */
417 template <typename Setter1Lambda, typename Setter2Lambda, typename UnselectedSetterLambda,
418 typename Select1Lambda, typename Select2Lambda>
439
440 /*!
441 * \brief Effectue un partitionnement d'une liste en trois parties.
442 *
443 * Le nombre de valeurs de la liste est donné par \a nb_value.
444 * Les deux fonctions lambda \a select1_lambda et \a select2_lambda permettent
445 * de partitionner la liste avec l'algorithme suivant:
446 * - si select1_lambda() est vrai, la valeur ajoutée \a first_output_iter,
447 * - sinon si select2_lambda() est vrai, la valeur sera ajoutée à \a second_output_iter,
448 * - sinon la valeur sera ajoutée à \a unselected_iter.
449 *
450 * Les listes en sortie sont dans le même ordre qu'en entrée.
451 *
452 * Après exécution, il est possible de récupérer le nombre d'éléments
453 * de la première et de la deuxième liste la méthode \a nbParts().
454 *
455 * \snippet AcceleratorPartitionerUnitTest.cc SampleListPartitionerThreeWayIf
456 */
457 template <typename InputIterator, typename FirstOutputIterator,
458 typename SecondOutputIterator, typename UnselectedIterator,
459 typename Select1Lambda, typename Select2Lambda>
476
477 /*!
478 * \brief Nombre d'éléments de la première partie de la liste.
479 */
481 {
482 m_is_already_called = false;
483 return _nbFirstPart();
484 }
485
486 /*!
487 * \brief Nombre d'éléments de la première et deuxième partie de la liste.
488 *
489 * Retourne une vue de deux valeurs. La première valeur contient le nombre
490 * d'éléments de la première liste et la seconde valeur le
491 * nombre d'éléments de la deuxième liste.
492 *
493 * Cette méthode n'est valide qu'après avoir appelé une méthode de partitionnement
494 * en trois parties.
495 */
497 {
498 m_is_already_called = false;
499 return _nbParts();
500 }
501
502 private:
503
504 bool m_is_already_called = false;
505
506 private:
507
508 void _setCalled()
509 {
510 if (m_is_already_called)
511 ARCANE_FATAL("apply() has already been called for this instance");
512 m_is_already_called = true;
513 }
514 bool _checkEmpty(Int32 nb_value)
515 {
516 if (nb_value == 0) {
517 _resetNbPart();
518 return true;
519 }
520 return false;
521 }
522};
523
524/*---------------------------------------------------------------------------*/
525/*---------------------------------------------------------------------------*/
526
527} // namespace Arcane::Accelerator
528
529/*---------------------------------------------------------------------------*/
530/*---------------------------------------------------------------------------*/
531
532#endif
533
534/*---------------------------------------------------------------------------*/
535/*---------------------------------------------------------------------------*/
#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.
Algorithme générique de partitionnement d'une liste.
SmallSpan< const Int32 > nbParts()
Nombre d'éléments de la première et deuxième partie de la liste.
Int32 nbFirstPart()
Nombre d'éléments de la première partie de la liste.
void applyWithIndex(Int32 nb_value, const Setter1Lambda setter1_lambda, const Setter2Lambda setter2_lambda, const UnselectedSetterLambda &unselected_setter_lambda, const Select1Lambda &select1_lambda, const Select2Lambda &select2_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue un partitionnement d'une liste en trois parties.
void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue un partitionnement d'une liste en deux parties.
void applyWithIndex(Int32 nb_value, const SetterLambda &setter_lambda, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue un partitionnement d'une liste en deux parties.
void applyIf(Int32 nb_value, InputIterator input_iter, FirstOutputIterator first_output_iter, SecondOutputIterator second_output_iter, UnselectedIterator unselected_iter, const Select1Lambda &select1_lambda, const Select2Lambda &select2_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue un partitionnement d'une liste en trois parties.
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
void copyToAsync(SmallSpan< DataType > dest_ptr, const RunQueue &queue)
Copie l'instance dans dest_ptr.
void apply3(GenericPartitionerBase &s, Int32 nb_item, InputIterator input_iter, FirstOutputIterator first_output_iter, SecondOutputIterator second_output_iter, UnselectedIterator unselected_iter, const Select1Lambda &select1_lambda, const Select2Lambda &select2_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue le partitionnement d'une liste en trois parties.
void apply(GenericPartitionerBase &s, Int32 nb_item, InputIterator input_iter, OutputIterator output_iter, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Effectue le partitionnement d'une liste en deux parties.
Itérateur sur une lambda pour positionner une valeur via un index.
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.