Arcane  v3.16.0.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>
80 void apply(GenericPartitionerBase& s, Int32 nb_item, InputIterator input_iter, OutputIterator output_iter,
81 const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
82 {
83 RunQueue queue = s.m_queue;
84 eExecutionPolicy exec_policy = queue.executionPolicy();
85 RunCommand command = makeCommand(queue);
86 command << trace_info;
87 impl::RunCommandLaunchInfo launch_info(command, nb_item);
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,
97 input_iter, output_iter, nb_list1_ptr, nb_item,
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,
103 input_iter, output_iter, nb_list1_ptr, nb_item,
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,
115 nb_list1_ptr, nb_item, select_lambda, stream));
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,
121 nb_list1_ptr, nb_item, select_lambda, stream));
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;
139 NumArray<DataType, MDDim1> tmp_output_numarray(nb_item);
140 NumArray<bool, MDDim1> tmp_select_numarray(nb_item);
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];
178 output_iter[i] = tmp_output[reverse_i];
179 output_iter[reverse_i] = tmp_output[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]];
192 auto saved_output_iter = output_iter;
193 auto output2_iter = output_iter + nb_item;
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>
222 InputIterator input_iter,
223 FirstOutputIterator first_output_iter,
224 SecondOutputIterator second_output_iter,
225 UnselectedIterator unselected_iter,
226 const Select1Lambda& select1_lambda,
227 const Select2Lambda& select2_lambda,
228 const TraceInfo& trace_info = TraceInfo())
229 {
230 RunQueue queue = s.m_queue;
231 eExecutionPolicy exec_policy = queue.executionPolicy();
232 RunCommand command = makeCommand(queue);
233 command << trace_info;
234 impl::RunCommandLaunchInfo launch_info(command, nb_item);
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,
244 input_iter, first_output_iter, second_output_iter,
245 unselected_iter, nb_list1_ptr, nb_item,
246 select1_lambda, select2_lambda, stream));
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,
251 input_iter, first_output_iter, second_output_iter,
252 unselected_iter, nb_list1_ptr, nb_item,
253 select1_lambda, select2_lambda, stream));
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,
265 second_output_iter, unselected_iter,
266 nb_list1_ptr, nb_item, select1_lambda, select2_lambda, stream));
267
268 s.m_algo_storage.allocate(temp_storage_size);
269 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
270
271 ARCANE_CHECK_HIP(partition_three_way(s.m_algo_storage.address(), temp_storage_size, input_iter, first_output_iter,
272 second_output_iter, unselected_iter, nb_list1_ptr, nb_item,
273 select1_lambda, select2_lambda, stream));
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) {
288 *first_output_iter = v;
289 ++first_output_iter;
290 ++nb_first;
291 }
292 else {
293 if (is_2) {
294 *second_output_iter = v;
295 ++second_output_iter;
296 ++nb_second;
297 }
298 else {
299 *unselected_iter = v;
300 ++unselected_iter;
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 */
332class GenericPartitioner
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>
358 void applyWithIndex(Int32 nb_value, const SetterLambda& setter_lambda,
359 const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
360 {
361 if (_checkEmpty(nb_value))
362 return;
363 _setCalled();
364 impl::GenericPartitionerBase* base_ptr = this;
366
367 impl::IndexIterator input_iter;
369 gf.apply(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
370 }
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>
389 void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter,
390 const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
391 {
392 if (_checkEmpty(nb_value))
393 return;
394 _setCalled();
395 impl::GenericPartitionerBase* base_ptr = this;
397 gf.apply(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
398 }
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>
419 void applyWithIndex(Int32 nb_value,
420 const Setter1Lambda setter1_lambda,
421 const Setter2Lambda setter2_lambda,
422 const UnselectedSetterLambda& unselected_setter_lambda,
423 const Select1Lambda& select1_lambda,
424 const Select2Lambda& select2_lambda,
425 const TraceInfo& trace_info = TraceInfo())
426 {
427 if (_checkEmpty(nb_value))
428 return;
429 _setCalled();
430 impl::GenericPartitionerBase* base_ptr = this;
432 impl::IndexIterator input_iter;
433 impl::SetterLambdaIterator<Setter1Lambda> setter1_wrapper(setter1_lambda);
434 impl::SetterLambdaIterator<Setter2Lambda> setter2_wrapper(setter2_lambda);
435 impl::SetterLambdaIterator<UnselectedSetterLambda> unselected_setter_wrapper(unselected_setter_lambda);
436 gf.apply3(*base_ptr, nb_value, input_iter, setter1_wrapper, setter2_wrapper,
437 unselected_setter_wrapper, select1_lambda, select2_lambda, trace_info);
438 }
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>
460 void applyIf(Int32 nb_value, InputIterator input_iter,
461 FirstOutputIterator first_output_iter,
462 SecondOutputIterator second_output_iter,
463 UnselectedIterator unselected_iter,
464 const Select1Lambda& select1_lambda,
465 const Select2Lambda& select2_lambda,
466 const TraceInfo& trace_info = TraceInfo())
467 {
468 if (_checkEmpty(nb_value))
469 return;
470 _setCalled();
471 impl::GenericPartitionerBase* base_ptr = this;
473 gf.apply3(*base_ptr, nb_value, input_iter, first_output_iter, second_output_iter,
474 unselected_iter, select1_lambda, select2_lambda, trace_info);
475 }
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.
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.
void beginExecute()
Indique qu'on commence l'exécution de la commande.
void endExecute()
Signale la fin de l'exécution.
Itérateur sur une lambda pour positionner une valeur via un index.
Tableaux multi-dimensionnels pour les types numériques accessibles sur accélérateurs.
constexpr SmallSpan< DataType > to1DSmallSpan()
Vue 1D sur l'instance (uniquement si rank == 1)
Vue d'un tableau d'éléments de type T.
Definition Span.h:673
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.
std::int32_t Int32
Type entier signé sur 32 bits.