12#ifndef ARCANE_ACCELERATOR_GENERICPARTITIONER_H
13#define ARCANE_ACCELERATOR_GENERICPARTITIONER_H
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
19#include "arcane/utils/NumArray.h"
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)
32namespace Arcane::Accelerator::impl
43class ARCANE_ACCELERATOR_EXPORT GenericPartitionerBase
45 friend class GenericPartitionerIf;
49 explicit GenericPartitionerBase(
const RunQueue& queue);
53 Int32 _nbFirstPart()
const;
79 template <
typename SelectLambda,
typename InputIterator,
typename OutputIterator>
86 command << trace_info;
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);
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));
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);
108#if defined(ARCANE_COMPILING_HIP)
110 size_t temp_storage_size = 0;
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));
117 s.m_algo_storage.allocate(temp_storage_size);
118 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
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);
125#if defined(ARCANE_COMPILING_SYCL) && defined(__INTEL_LLVM_COMPILER)
137 using InputDataType =
typename InputIterator::value_type;
138 using DataType =
typename OutputIterator::value_type;
148 tmp_output[i] = input_iter[i];
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);
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;
167 Int32 nb_iter2 = (nb_list2 / 2) + (nb_list2 % 2);
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];
182 output_iter[i] = tmp_output[i];
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)) {
206 Int32 nb_list1 =
static_cast<Int32>(output_iter - saved_output_iter);
207 s.m_host_nb_list1_storage[0] = nb_list1;
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,
233 command << trace_info;
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);
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));
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);
257#if defined(ARCANE_COMPILING_HIP)
259 size_t temp_storage_size = 0;
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));
268 s.m_algo_storage.allocate(temp_storage_size);
269 nb_list1_ptr = s.m_device_nb_list1_storage.allocate();
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);
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);
288 *first_output_iter = v;
294 *second_output_iter = v;
295 ++second_output_iter;
299 *unselected_iter = v;
306 s.m_host_nb_list1_storage[0] = nb_first;
307 s.m_host_nb_list1_storage[1] = nb_second;
332class GenericPartitioner
337 explicit GenericPartitioner(
const RunQueue& queue)
357 template <
typename SelectLambda,
typename SetterLambda>
361 if (_checkEmpty(nb_value))
369 gf.
apply(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
388 template <
typename InputIterator,
typename OutputIterator,
typename SelectLambda>
389 void applyIf(
Int32 nb_value, InputIterator input_iter, OutputIterator output_iter,
392 if (_checkEmpty(nb_value))
397 gf.
apply(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
417 template <
typename Setter1Lambda,
typename Setter2Lambda,
typename UnselectedSetterLambda,
418 typename Select1Lambda,
typename Select2Lambda>
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,
427 if (_checkEmpty(nb_value))
436 gf.
apply3(*base_ptr, nb_value, input_iter, setter1_wrapper, setter2_wrapper,
437 unselected_setter_wrapper, select1_lambda, select2_lambda, trace_info);
457 template <
typename InputIterator,
typename FirstOutputIterator,
458 typename SecondOutputIterator,
typename UnselectedIterator,
459 typename Select1Lambda,
typename Select2Lambda>
461 FirstOutputIterator first_output_iter,
462 SecondOutputIterator second_output_iter,
463 UnselectedIterator unselected_iter,
464 const Select1Lambda& select1_lambda,
465 const Select2Lambda& select2_lambda,
468 if (_checkEmpty(nb_value))
473 gf.
apply3(*base_ptr, nb_value, input_iter, first_output_iter, second_output_iter,
474 unselected_iter, select1_lambda, select2_lambda, trace_info);
482 m_is_already_called =
false;
483 return _nbFirstPart();
498 m_is_already_called =
false;
504 bool m_is_already_called =
false;
510 if (m_is_already_called)
511 ARCANE_FATAL(
"apply() has already been called for this instance");
512 m_is_already_called =
true;
514 bool _checkEmpty(
Int32 nb_value)
#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.
eExecutionPolicy executionPolicy() const
Politique d'exécution de la file.
Gère l'allocation interne sur le device pour un type donné.
void copyToAsync(SmallSpan< DataType > dest_ptr, const RunQueue &queue)
Copie l'instance dans dest_ptr.
Gère l'allocation interne sur le device.
Classe de base pour effectuer un filtrage.
Classe pour effectuer un partitionnement d'une liste.
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.
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.
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.
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.