12#ifndef ARCANE_ACCELERATOR_FILTERING_H
13#define ARCANE_ACCELERATOR_FILTERING_H
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
19#include "arcane/utils/NumArray.h"
20#include "arcane/utils/TraceInfo.h"
22#include "arcane/accelerator/core/RunQueue.h"
24#include "arcane/accelerator/AcceleratorGlobal.h"
25#include "arcane/accelerator/CommonUtils.h"
26#include "arcane/accelerator/RunCommandLaunchInfo.h"
28#include "arcane/accelerator/ScanImpl.h"
33namespace Arcane::Accelerator::impl
47 template <
typename DataType,
typename FlagType,
typename OutputDataType>
50 friend class SyclGenericFilteringImpl;
61 Int32 _nbOutputElement()
const;
63 void _allocateTemporaryStorage(
size_t size);
64 int* _getDeviceNbOutPointer();
65 void _copyDeviceNbOutToHostNbOut();
84 bool m_use_direct_host_storage =
true;
90#if defined(ARCANE_COMPILING_SYCL)
92class SyclGenericFilteringImpl
96 template <
typename SelectLambda,
typename InputIterator,
typename OutputIterator>
98 OutputIterator output_iter, SelectLambda select_lambda)
101 using DataType = std::iterator_traits<OutputIterator>::value_type;
102#if defined(ARCANE_USE_SCAN_ONEDPL) && defined(__INTEL_LLVM_COMPILER)
103 sycl::queue true_queue = AcceleratorUtils::toSyclNativeStream(queue);
104 auto policy = oneapi::dpl::execution::make_device_policy(true_queue);
105 auto out_iter = oneapi::dpl::copy_if(policy, input_iter, input_iter + nb_item, output_iter, select_lambda);
106 Int32 nb_output = out_iter - output_iter;
118 in_scan_data[i] = select_lambda(input_iter[i]) ? 1 : 0;
122 SyclScanner<
false ,
Int32, ScannerSumOperator<Int32>> scanner;
123 scanner.doScan(queue, in_scan_data, out_scan_data, 0);
125 Int32 nb_output = out_scan_data[nb_item - 1];
128 const bool do_verbose =
false;
129 if (do_verbose && nb_item < 1500)
130 for (
int i = 0; i < nb_item; ++i) {
131 std::cout <<
"out_data i=" << i <<
" out_data=" << out_scan_data[i]
132 <<
" in_data=" << in_scan_data[i] <<
" value=" << input_iter[i] <<
"\n ";
139 out_copy.
resize(nb_output);
140 auto out_copy_view = out_copy.
to1DSpan();
146 if (in_scan_data[i] == 1)
147 out_copy_view[out_scan_data[i] - 1] = input_iter[i];
155 output_iter[i] = out_copy_view[i];
175template <
typename DataType,
typename FlagType,
typename OutputDataType>
183 if (output.
size() != nb_item)
184 ARCANE_FATAL(
"Sizes are not equals: input={0} output={1}", nb_item, output.
size());
185 [[maybe_unused]]
const DataType* input_data = input.
data();
186 [[maybe_unused]] DataType* output_data = output.
data();
187 [[maybe_unused]]
const FlagType* flag_data = flag.
data();
191 switch (exec_policy) {
192#if defined(ARCANE_COMPILING_CUDA)
194 size_t temp_storage_size = 0;
195 cudaStream_t stream = AcceleratorUtils::toCudaNativeStream(queue);
197 int* nb_out_ptr =
nullptr;
198 ARCANE_CHECK_CUDA(::cub::DeviceSelect::Flagged(
nullptr, temp_storage_size,
199 input_data, flag_data, output_data, nb_out_ptr, nb_item, stream));
201 s._allocateTemporaryStorage(temp_storage_size);
202 nb_out_ptr = s._getDeviceNbOutPointer();
203 ARCANE_CHECK_CUDA(::cub::DeviceSelect::Flagged(s.m_algo_storage.address(), temp_storage_size,
204 input_data, flag_data, output_data, nb_out_ptr, nb_item, stream));
205 s._copyDeviceNbOutToHostNbOut();
208#if defined(ARCANE_COMPILING_HIP)
210 size_t temp_storage_size = 0;
212 hipStream_t stream = AcceleratorUtils::toHipNativeStream(queue);
213 int* nb_out_ptr =
nullptr;
214 ARCANE_CHECK_HIP(rocprim::select(
nullptr, temp_storage_size, input_data, flag_data, output_data,
215 nb_out_ptr, nb_item, stream));
217 s._allocateTemporaryStorage(temp_storage_size);
218 nb_out_ptr = s._getDeviceNbOutPointer();
220 ARCANE_CHECK_HIP(rocprim::select(s.m_algo_storage.address(), temp_storage_size, input_data, flag_data, output_data,
221 nb_out_ptr, nb_item, stream));
222 s._copyDeviceNbOutToHostNbOut();
225#if defined(ARCANE_COMPILING_SYCL)
228 auto filter_lambda = [=](
Int32 input_index) ->
bool {
return flag[input_index] != 0; };
229 auto setter_lambda = [=](
Int32 input_index,
Int32 output_index) { output[output_index] = input[input_index]; };
231 SyclGenericFilteringImpl::apply(s, nb_item, iter2, out, filter_lambda);
239 for (
Int32 i = 0; i < nb_item; ++i) {
241 output[index] = input[i];
266 template <
typename SelectLambda,
typename InputIterator,
typename OutputIterator>
268 const SelectLambda& select_lambda,
const TraceInfo& trace_info)
274 command << trace_info;
277 switch (exec_policy) {
278#if defined(ARCANE_COMPILING_CUDA)
280 size_t temp_storage_size = 0;
281 cudaStream_t stream = impl::CudaUtils::toNativeStream(queue);
283 int* nb_out_ptr =
nullptr;
284 ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(
nullptr, temp_storage_size,
285 input_iter, output_iter, nb_out_ptr, nb_item,
286 select_lambda, stream));
288 s._allocateTemporaryStorage(temp_storage_size);
289 nb_out_ptr = s._getDeviceNbOutPointer();
290 ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(s.m_algo_storage.address(), temp_storage_size,
291 input_iter, output_iter, nb_out_ptr, nb_item,
292 select_lambda, stream));
293 s._copyDeviceNbOutToHostNbOut();
296#if defined(ARCANE_COMPILING_HIP)
298 size_t temp_storage_size = 0;
300 hipStream_t stream = impl::HipUtils::toNativeStream(queue);
301 int* nb_out_ptr =
nullptr;
302 ARCANE_CHECK_HIP(rocprim::select(
nullptr, temp_storage_size, input_iter, output_iter,
303 nb_out_ptr, nb_item, select_lambda, stream));
305 s._allocateTemporaryStorage(temp_storage_size);
306 nb_out_ptr = s._getDeviceNbOutPointer();
307 ARCANE_CHECK_HIP(rocprim::select(s.m_algo_storage.address(), temp_storage_size, input_iter, output_iter,
308 nb_out_ptr, nb_item, select_lambda, 0));
309 s._copyDeviceNbOutToHostNbOut();
312#if defined(ARCANE_COMPILING_SYCL)
314 SyclGenericFilteringImpl::apply(s, nb_item, input_iter, output_iter, select_lambda);
322 for (
Int32 i = 0; i < nb_item; ++i) {
323 if (select_lambda(*input_iter)) {
324 *output_iter = *input_iter;
363 ARCANE_DEPRECATED_REASON(
"Y2024: Use GenericFilterer(const RunQueue&) instead")
407 template <
typename InputDataType,
typename OutputDataType,
typename FlagType>
410 const Int32 nb_value = input.
size();
411 if (output.
size() != nb_value)
412 ARCANE_FATAL(
"Sizes are not equals: input={0} output={1}", nb_value, output.
size());
413 if (flag.
size() != nb_value)
414 ARCANE_FATAL(
"Sizes are not equals: input={0} flag={1}", nb_value, flag.
size());
417 if (_checkEmpty(nb_value))
421 gf.apply(*base_ptr, input, output, flag);
458 template <
typename DataType,
typename SelectLambda>
462 const Int32 nb_value = input.
size();
463 if (output.
size() != nb_value)
464 ARCANE_FATAL(
"Sizes are not equals: input={0} output={1}", nb_value, output.
size());
467 if (_checkEmpty(nb_value))
471 gf.apply(*base_ptr, nb_value, input.
data(), output.
data(), select_lambda, trace_info);
482 template <
typename InputIterator,
typename OutputIterator,
typename SelectLambda>
483 void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter,
487 if (_checkEmpty(nb_value))
491 gf.apply(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
525 template <
typename SelectLambda,
typename SetterLambda>
530 if (_checkEmpty(nb_value))
536 gf.apply(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
542 m_is_already_called =
false;
543 return _nbOutputElement();
548 bool m_is_already_called =
false;
554 if (m_is_already_called)
555 ARCANE_FATAL(
"apply() has already been called for this instance");
556 m_is_already_called =
true;
558 bool _checkEmpty(Int32 nb_value)
#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.
Algorithme générique de filtrage sur accélérateur.
void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique un filtre.
void applyWithIndex(Int32 nb_value, const SelectLambda &select_lambda, const SetterLambda &setter_lambda, const TraceInfo &trace_info=TraceInfo())
Applique un filtre avec une sélection suivant un index.
void apply(SmallSpan< const InputDataType > input, SmallSpan< OutputDataType > output, SmallSpan< const FlagType > flag)
Applique un filtre.
void applyIf(SmallSpan< const DataType > input, SmallSpan< DataType > output, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique un filtre.
Int32 nbOutputElement()
Nombre d'éléments en sortie.
GenericFilterer(const RunQueue &queue)
Créé une instance.
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.
DeviceStorage< int > m_device_nb_out_storage
Mémoire sur le device du nombre de valeurs filtrées.
NumArray< Int32, MDDim1 > m_host_nb_out_storage
Mémoire hôte pour le nombre de valeurs filtrées.
RunQueue m_queue
File d'exécution. Ne doit pas être nulle.
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.
Span< const DataType > to1DSpan() const
Vue 1D constante sur l'instance.
constexpr SmallSpan< DataType > to1DSmallSpan()
Vue 1D sur l'instance (uniquement si rank == 1)
void resize(Int32 dim1_size, Int32 dim2_size, Int32 dim3_size, Int32 dim4_size)
Modifie la taille du tableau.
Vue d'un tableau d'éléments de type T.
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
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.