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
53 Int32 _nbFirstPart()
const;
79 template <
typename SelectLambda,
typename InputIterator,
typename OutputIterator>
90#if defined(ARCANE_COMPILING_CUDA)
93 cudaStream_t stream = impl::CudaUtils::toNativeStream(&queue);
105 s.m_device_nb_list1_storage.
copyToAsync(s.m_host_nb_list1_storage, queue);
108#if defined(ARCANE_COMPILING_HIP)
112 hipStream_t stream = impl::HipUtils::toNativeStream(&queue);
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)
138 using DataType =
typename OutputIterator::value_type;
152 sycl::queue
sycl_queue = impl::SyclUtils::toNativeStream(queue);
158 s.m_host_nb_list1_storage[0] =
nb_list1;
194 for (Int32 i = 0; i <
nb_item; ++i) {
207 s.m_host_nb_list1_storage[0] =
nb_list1;
237#if defined(ARCANE_COMPILING_CUDA)
240 cudaStream_t stream = impl::CudaUtils::toNativeStream(&queue);
254 s.m_device_nb_list1_storage.
copyToAsync(s.m_host_nb_list1_storage, queue);
257#if defined(ARCANE_COMPILING_HIP)
261 hipStream_t stream = impl::HipUtils::toNativeStream(&queue);
274 s.m_device_nb_list1_storage.
copyToAsync(s.m_host_nb_list1_storage, queue);
283 for (Int32 i = 0; i <
nb_item; ++i) {
306 s.m_host_nb_list1_storage[0] =
nb_first;
307 s.m_host_nb_list1_storage[1] =
nb_second;
357 template <
typename SelectLambda,
typename SetterLambda>
388 template <
typename InputIterator,
typename OutputIterator,
typename SelectLambda>
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;
#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.
eExecutionPolicy executionPolicy() const
Politique d'exécution de la file.
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.