12#ifndef ARCANE_ACCELERATOR_GENERICSCANNER_H
13#define ARCANE_ACCELERATOR_GENERICSCANNER_H
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
20#include "arcane/utils/NumArray.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"
29#include "arcane/accelerator/MultiThreadAlgo.h"
34namespace Arcane::Accelerator::impl
54 typename Operator,
typename DataType>
64#if defined(ARCANE_COMPILING_CUDA)
67 cudaStream_t stream = impl::CudaUtils::toNativeStream(&m_queue);
84#if defined(ARCANE_COMPILING_HIP)
88 hipStream_t stream = impl::HipUtils::toNativeStream(&m_queue);
104#if defined(ARCANE_COMPILING_SYCL)
106#if defined(ARCANE_USE_SCAN_ONEDPL) && defined(__INTEL_LLVM_COMPILER)
107 sycl::queue queue = impl::SyclUtils::toNativeStream(&m_queue);
108 auto policy = oneapi::dpl::execution::make_device_policy(queue);
156 for (Int32 i = 0; i <
nb_item; ++i) {
199template <
typename DataType>
210 static void exclusiveMin(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
212 _applyArray<true>(queue, input, output, ScannerMinOperator<DataType>{});
215 static void exclusiveMax(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
217 _applyArray<true>(queue, input, output, ScannerMaxOperator<DataType>{});
220 static void inclusiveSum(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
222 _applyArray<false>(queue, input, output, ScannerSumOperator<DataType>{});
225 static void inclusiveMin(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
227 _applyArray<false>(queue, input, output, ScannerMinOperator<DataType>{});
230 static void inclusiveMax(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output)
232 _applyArray<false>(queue, input, output, ScannerMaxOperator<DataType>{});
237 template <
bool IsExclusive,
typename Operator>
238 static void _applyArray(RunQueue* queue, SmallSpan<const DataType> input, SmallSpan<DataType> output,
const Operator& op)
241 impl::ScannerImpl scanner(*queue);
242 const Int32 nb_item = input.size();
243 if (output.size() != nb_item)
244 ARCANE_FATAL(
"Sizes are not equals: input={0} output={1}", nb_item, output.size());
245 const DataType* input_data = input.data();
246 DataType* output_data = output.data();
247 DataType init_value = op.defaultValue();
248 scanner.apply<IsExclusive>(nb_item, input_data, output_data, init_value, op, TraceInfo{});
269 template <
typename DataType,
typename SetterLambda>
285 m_lambda(m_index, value);
294 using value_type = DataType;
295 using iterator_category = std::random_access_iterator_tag;
320 return ThatClass(iter.m_lambda, iter.m_index + x);
322 ARCCORE_HOST_DEVICE
friend ThatClass operator+(
Int32 x,
const ThatClass& iter)
324 return ThatClass(iter.m_lambda, iter.m_index + x);
326 ARCCORE_HOST_DEVICE
friend bool operator<(
const ThatClass& iter1,
const ThatClass& iter2)
328 return iter1.m_index < iter2.m_index;
330 ARCCORE_HOST_DEVICE ThatClass operator-(
Int32 x)
332 return ThatClass(m_lambda, m_index - x);
334 ARCCORE_HOST_DEVICE
Int32 operator-(
const ThatClass& x)
const
336 return m_index - x.m_index;
338 ARCCORE_HOST_DEVICE reference operator*()
const
340 return Setter(m_lambda, m_index);
342 ARCCORE_HOST_DEVICE reference operator[](
Int32 x)
const {
return Setter(m_lambda, m_index + x); }
343 ARCCORE_HOST_DEVICE
friend bool operator!=(
const ThatClass& a,
const ThatClass& b)
345 return a.m_index != b.m_index;
351 SetterLambda m_lambda;
356 explicit GenericScanner(
const RunQueue& queue)
362 template <
typename DataType,
typename GetterLambda,
typename SetterLambda,
typename Operator>
363 void applyWithIndexExclusive(
Int32 nb_value,
const DataType& initial_value,
364 const GetterLambda& getter_lambda,
365 const SetterLambda& setter_lambda,
366 const Operator& op_lambda,
367 const TraceInfo& trace_info = TraceInfo())
369 _applyWithIndex<true>(nb_value, initial_value, getter_lambda, setter_lambda, op_lambda, trace_info);
372 template <
typename DataType,
typename GetterLambda,
typename SetterLambda,
typename Operator>
373 void applyWithIndexInclusive(
Int32 nb_value,
const DataType& initial_value,
374 const GetterLambda& getter_lambda,
375 const SetterLambda& setter_lambda,
376 const Operator& op_lambda,
377 const TraceInfo& trace_info = TraceInfo())
379 _applyWithIndex<false>(nb_value, initial_value, getter_lambda, setter_lambda, op_lambda, trace_info);
382 template <
typename InputDataType,
typename OutputDataType,
typename Operator>
383 void applyExclusive(
const OutputDataType& initial_value,
384 SmallSpan<const InputDataType> input,
385 SmallSpan<OutputDataType> output,
386 const Operator& op_lambda,
387 const TraceInfo& trace_info = TraceInfo())
389 _apply<true>(initial_value, input, output, op_lambda, trace_info);
392 template <
typename InputDataType,
typename OutputDataType,
typename Operator>
393 void applyInclusive(
const OutputDataType& initial_value,
394 SmallSpan<const InputDataType> input,
395 SmallSpan<OutputDataType> output,
396 const Operator& op_lambda,
397 const TraceInfo& trace_info = TraceInfo())
399 _apply<false>(initial_value, input, output, op_lambda, trace_info);
404 template <
bool IsExclusive,
typename DataType,
typename GetterLambda,
typename SetterLambda,
typename Operator>
405 void _applyWithIndex(
Int32 nb_value,
const DataType& initial_value,
406 const GetterLambda& getter_lambda,
407 const SetterLambda& setter_lambda,
408 const Operator& op_lambda,
409 const TraceInfo& trace_info)
411 impl::GetterLambdaIterator<DataType, GetterLambda> input_iter(getter_lambda);
412 SetterLambdaIterator<DataType, SetterLambda> output_iter(setter_lambda);
413 impl::ScannerImpl scanner(m_queue);
414 scanner.apply<IsExclusive>(nb_value, input_iter, output_iter, initial_value, op_lambda, trace_info);
417 template <
bool IsExclusive,
typename InputDataType,
typename OutputDataType,
typename Operator>
418 void _apply(
const OutputDataType& initial_value,
419 SmallSpan<const InputDataType> input,
420 SmallSpan<OutputDataType> output,
422 const TraceInfo& trace_info = TraceInfo())
424 const Int32 nb_item = input.size();
425 if (output.size() != nb_item)
426 ARCANE_FATAL(
"Sizes are not equals: input={0} output={1}", nb_item, output.size());
427 auto* input_data = input.data();
428 auto* output_data = output.data();
429 impl::ScannerImpl scanner(m_queue);
430 scanner.apply<IsExclusive>(nb_item, input_data, output_data, initial_value, op, trace_info);
#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.
Permet de positionner un élément de l'itérateur de sortie.
Itérateur sur une lambda pour positionner une valeur via un index.
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.
Algorithmes avancée en mode multi-thread.
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.