Arcane  v3.14.10.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
Filter.h
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2024 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/* Filtering.h (C) 2000-2024 */
9/* */
10/* Algorithme de filtrage. */
11/*---------------------------------------------------------------------------*/
12#ifndef ARCANE_ACCELERATOR_FILTERING_H
13#define ARCANE_ACCELERATOR_FILTERING_H
14/*---------------------------------------------------------------------------*/
15/*---------------------------------------------------------------------------*/
16
17#include "arcane/utils/ArrayView.h"
18#include "arcane/utils/FatalErrorException.h"
19#include "arcane/utils/NumArray.h"
20#include "arcane/utils/TraceInfo.h"
21
22#include "arcane/accelerator/core/RunQueue.h"
23
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
30/*---------------------------------------------------------------------------*/
31/*---------------------------------------------------------------------------*/
32
33namespace Arcane::Accelerator::impl
34{
35//#define ARCANE_USE_SCAN_ONEDPL
36
37/*---------------------------------------------------------------------------*/
38/*---------------------------------------------------------------------------*/
39/*!
40 * \internal
41 * \brief Classe de base pour effectuer un filtrage.
42 *
43 * Contient les arguments nécessaires pour effectuer le filtrage.
44 */
45class ARCANE_ACCELERATOR_EXPORT GenericFilteringBase
46{
47 template <typename DataType, typename FlagType, typename OutputDataType>
48 friend class GenericFilteringFlag;
49 friend class GenericFilteringIf;
50 friend class SyclGenericFilteringImpl;
51
52 public:
53
54
55 protected:
56
58
59 protected:
60
61 Int32 _nbOutputElement() const;
62 void _allocate();
63 void _allocateTemporaryStorage(size_t size);
64 int* _getDeviceNbOutPointer();
65 void _copyDeviceNbOutToHostNbOut();
66
67 protected:
68
69 //! File d'exécution. Ne doit pas être nulle.
71 // Mémoire de travail pour l'algorithme de filtrage.
72 GenericDeviceStorage m_algo_storage;
73 //! Mémoire sur le device du nombre de valeurs filtrées
75 //! Mémoire hôte pour le nombre de valeurs filtrées
77 /*!
78 * \brief Indique quelle mémoire est utilisée pour le nombre de valeurs filtrées.
79 *
80 * Si vrai utilise directement \a m_host_nb_out_storage. Sinon, utilise
81 * m_device_nb_out_storage et fait une copie asynchrone après le filtrage pour
82 * recopier la valeur dans m_host_nb_out_storage.
83 */
84 bool m_use_direct_host_storage = true;
85};
86
87/*---------------------------------------------------------------------------*/
88/*---------------------------------------------------------------------------*/
89
90#if defined(ARCANE_COMPILING_SYCL)
91//! Implémentation pour SYCL
92class SyclGenericFilteringImpl
93{
94 public:
95
96 template <typename SelectLambda, typename InputIterator, typename OutputIterator>
97 static void apply(GenericFilteringBase& s, Int32 nb_item, InputIterator input_iter,
98 OutputIterator output_iter, SelectLambda select_lambda)
99 {
100 RunQueue queue = s.m_queue;
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;
107 s.m_host_nb_out_storage[0] = nb_output;
108#else
109 NumArray<Int32, MDDim1> scan_input_data(nb_item);
110 NumArray<Int32, MDDim1> scan_output_data(nb_item);
111 SmallSpan<Int32> in_scan_data = scan_input_data.to1DSmallSpan();
112 SmallSpan<Int32> out_scan_data = scan_output_data.to1DSmallSpan();
113 {
114 auto command = makeCommand(queue);
115 command << RUNCOMMAND_LOOP1(iter, nb_item)
116 {
117 auto [i] = iter();
118 in_scan_data[i] = select_lambda(input_iter[i]) ? 1 : 0;
119 };
120 }
121 queue.barrier();
122 SyclScanner<false /*is_exclusive*/, Int32, ScannerSumOperator<Int32>> scanner;
123 scanner.doScan(queue, in_scan_data, out_scan_data, 0);
124 // La valeur de 'out_data' pour le dernier élément (nb_item-1) contient la taille du filtre
125 Int32 nb_output = out_scan_data[nb_item - 1];
126 s.m_host_nb_out_storage[0] = nb_output;
127
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 ";
133 }
134 // Copie depuis 'out_data' vers 'in_data' les indices correspondant au filtre
135 // Comme 'output_iter' et 'input_iter' peuvent se chevaucher, il
136 // faut faire une copie intermédiaire
137 // TODO: détecter cela et ne faire la copie que si nécessaire.
138 NumArray<DataType,MDDim1> out_copy(eMemoryRessource::Device);
139 out_copy.resize(nb_output);
140 auto out_copy_view = out_copy.to1DSpan();
141 {
142 auto command = makeCommand(queue);
143 command << RUNCOMMAND_LOOP1(iter, nb_item)
144 {
145 auto [i] = iter();
146 if (in_scan_data[i] == 1)
147 out_copy_view[out_scan_data[i] - 1] = input_iter[i];
148 };
149 }
150 {
151 auto command = makeCommand(queue);
152 command << RUNCOMMAND_LOOP1(iter, nb_output)
153 {
154 auto [i] = iter();
155 output_iter[i] = out_copy_view[i];
156 };
157 }
158 // Obligatoire à cause de 'out_copy'. On pourra le supprimer avec une
159 // allocation temporaire.
160 queue.barrier();
161#endif
162 }
163};
164#endif
165
166/*---------------------------------------------------------------------------*/
167/*---------------------------------------------------------------------------*/
168/*!
169 * \internal
170 * \brief Classe pour effectuer un filtrage
171 *
172 * \a DataType est le type de donnée.
173 * \a FlagType est le type du tableau de filtre.
174 */
175template <typename DataType, typename FlagType, typename OutputDataType>
177{
178 public:
179
181 {
182 const Int32 nb_item = input.size();
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();
189 RunQueue queue = s.m_queue;
190 exec_policy = queue.executionPolicy();
191 switch (exec_policy) {
192#if defined(ARCANE_COMPILING_CUDA)
194 size_t temp_storage_size = 0;
195 cudaStream_t stream = AcceleratorUtils::toCudaNativeStream(queue);
196 // Premier appel pour connaitre la taille pour l'allocation
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));
200
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();
206 } break;
207#endif
208#if defined(ARCANE_COMPILING_HIP)
210 size_t temp_storage_size = 0;
211 // Premier appel pour connaitre la taille pour l'allocation
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));
216
217 s._allocateTemporaryStorage(temp_storage_size);
218 nb_out_ptr = s._getDeviceNbOutPointer();
219
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();
223 } break;
224#endif
225#if defined(ARCANE_COMPILING_SYCL)
227 impl::IndexIterator iter2(0);
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]; };
230 impl::SetterLambdaIterator<decltype(setter_lambda)> out(setter_lambda);
231 SyclGenericFilteringImpl::apply(s, nb_item, iter2, out, filter_lambda);
232 } break;
233#endif
235 // Pas encore implémenté en multi-thread
236 [[fallthrough]];
238 Int32 index = 0;
239 for (Int32 i = 0; i < nb_item; ++i) {
240 if (flag[i] != 0) {
241 output[index] = input[i];
242 ++index;
243 }
244 }
245 s.m_host_nb_out_storage[0] = index;
246 } break;
247 default:
248 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
249 }
250 }
251};
252
253/*---------------------------------------------------------------------------*/
254/*---------------------------------------------------------------------------*/
255/*!
256 * \internal
257 * \brief Classe pour effectuer un filtrage
258 *
259 * \a DataType est le type de donnée.
260 * \a FlagType est le type du tableau de filtre.
261 */
263{
264 public:
265
266 template <typename SelectLambda, typename InputIterator, typename OutputIterator>
267 void apply(GenericFilteringBase& s, Int32 nb_item, InputIterator input_iter, OutputIterator output_iter,
268 const SelectLambda& select_lambda, const TraceInfo& trace_info)
269 {
271 RunQueue queue = s.m_queue;
272 exec_policy = queue.executionPolicy();
273 RunCommand command = makeCommand(queue);
274 command << trace_info;
275 impl::RunCommandLaunchInfo launch_info(command, nb_item);
276 launch_info.beginExecute();
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);
282 // Premier appel pour connaitre la taille pour l'allocation
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));
287
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();
294 } break;
295#endif
296#if defined(ARCANE_COMPILING_HIP)
298 size_t temp_storage_size = 0;
299 // Premier appel pour connaitre la taille pour l'allocation
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));
304
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();
310 } break;
311#endif
312#if defined(ARCANE_COMPILING_SYCL)
314 SyclGenericFilteringImpl::apply(s, nb_item, input_iter, output_iter, select_lambda);
315 } break;
316#endif
318 // Pas encore implémenté en multi-thread
319 [[fallthrough]];
321 Int32 index = 0;
322 for (Int32 i = 0; i < nb_item; ++i) {
323 if (select_lambda(*input_iter)) {
324 *output_iter = *input_iter;
325 ++index;
326 ++output_iter;
327 }
328 ++input_iter;
329 }
330 s.m_host_nb_out_storage[0] = index;
331 } break;
332 default:
333 ARCANE_FATAL(getBadPolicyMessage(exec_policy));
334 }
335 launch_info.endExecute();
336 }
337};
338
339/*---------------------------------------------------------------------------*/
340/*---------------------------------------------------------------------------*/
341
342} // namespace Arcane::Accelerator::impl
343
344namespace Arcane::Accelerator
345{
346
347/*---------------------------------------------------------------------------*/
348/*---------------------------------------------------------------------------*/
349/*!
350 * \brief Algorithme générique de filtrage sur accélérateur.
351 */
354{
355
356 public:
357
358 /*!
359 * \brief Créé une instance.
360 *
361 * \pre queue!=nullptr
362 */
363 ARCANE_DEPRECATED_REASON("Y2024: Use GenericFilterer(const RunQueue&) instead")
364 explicit GenericFilterer(RunQueue* queue)
365 {
367 m_queue = *queue;
368 _allocate();
369 }
370
371 /*!
372 * \brief Créé une instance.
373 *
374 * \pre queue!=nullptr
375 */
376 explicit GenericFilterer(const RunQueue& queue)
377 {
378 m_queue = queue;
379 _allocate();
380 }
381
382 public:
383
384 /*!
385 * \brief Applique un filtre.
386 *
387 * Filtre tous les éléments de \a input pour lesquels \a flag est différent de 0 et
388 * remplit \a output avec les valeurs filtrées. \a output doit avoir une taille assez
389 * grande pour contenir tous les éléments filtrés.
390 *
391 * L'algorithme séquentiel est le suivant:
392 *
393 * \code
394 * Int32 index = 0;
395 * for (Int32 i = 0; i < nb_item; ++i) {
396 * if (flag[i] != 0) {
397 * output[index] = input[i];
398 * ++index;
399 * }
400 * }
401 * return index;
402 * \endcode
403 *
404 * Il faut appeler la méthode nbOutputElement() pour obtenir le nombre d'éléments
405 * après filtrage.
406 */
407 template <typename InputDataType, typename OutputDataType, typename FlagType>
409 {
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());
415
416 _setCalled();
417 if (_checkEmpty(nb_value))
418 return;
419 impl::GenericFilteringBase* base_ptr = this;
421 gf.apply(*base_ptr, input, output, flag);
422 }
423
424 /*!
425 * \brief Applique un filtre.
426 *
427 * Filtre tous les éléments de \a input pour lesquels \a select_lambda vaut \a true et
428 * remplit \a output avec les valeurs filtrées. \a output doit avoir une taille assez
429 * grande pour contenir tous les éléments filtrés.
430 *
431 * \a select_lambda doit avoir un opérateur `ARCCORE_HOST_DEVICE bool operator()(const DataType& v) const`.
432 *
433 * Par exemple la lambda suivante permet de ne garder que les éléments dont
434 * la valeur est supérieure à 569.
435 *
436 * \code
437 * auto filter_lambda = [] ARCCORE_HOST_DEVICE (const DataType& x) -> bool {
438 * return (x > 569.0);
439 * };
440 * \endcode
441 *
442 * L'algorithme séquentiel est le suivant:
443 *
444 * \code
445 * Int32 index = 0;
446 * for (Int32 i = 0; i < nb_item; ++i) {
447 * if (select_lambda(input[i])) {
448 * output[index] = input[i];
449 * ++index;
450 * }
451 * }
452 * return index;
453 * \endcode
454 *
455 * Il faut appeler la méthode nbOutputElement() pour obtenir le nombre d'éléments
456 * après filtrage.
457 */
458 template <typename DataType, typename SelectLambda>
460 const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
461 {
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());
465
466 _setCalled();
467 if (_checkEmpty(nb_value))
468 return;
469 impl::GenericFilteringBase* base_ptr = this;
471 gf.apply(*base_ptr, nb_value, input.data(), output.data(), select_lambda, trace_info);
472 }
473
474 /*!
475 * \brief Applique un filtre.
476 *
477 * Cette méthode est identique à Filterer::applyIf(SmallSpan<const DataType> input,
478 * SmallSpan<DataType> output, const SelectLambda& select_lambda) mais permet de spécifier un
479 * itérateur \a input_iter pour l'entrée et \a output_iter pour la sortie.
480 * Le nombre d'entité en entrée est donné par \a nb_value.
481 */
482 template <typename InputIterator, typename OutputIterator, typename SelectLambda>
483 void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter,
484 const SelectLambda& select_lambda, const TraceInfo& trace_info = TraceInfo())
485 {
486 _setCalled();
487 if (_checkEmpty(nb_value))
488 return;
489 impl::GenericFilteringBase* base_ptr = this;
491 gf.apply(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
492 }
493
494 /*!
495 * \brief Applique un filtre avec une sélection suivant un index.
496 *
497 * Cette méthode permet de filtrer en spécifiant une fonction lambda à la fois
498 * pour la sélection et l'affection. Le prototype de ces lambda est:
499 *
500 * \code
501 * auto select_lambda = [=] ARCCORE_HOST_DEVICE (Int32 index) -> bool;
502 * auto setter_lambda = [=] ARCCORE_HOST_DEVICE (Int32 input_index,Int32 output_index) -> void;
503 * \endcode
504 *
505 * Par exemple, si on souhaite recopier dans \a output le tableau \a input dont les
506 * valeurs sont supérieures à 25.0.
507 *
508 * \code
509 * SmallSpan<Real> input = ...;
510 * SmallSpan<Real> output = ...;
511 * auto select_lambda = [=] ARCCORE_HOST_DEVICE (Int32 index) -> bool
512 * {
513 * return input[index] > 25.0;
514 * };
515 * auto setter_lambda = [=] ARCCORE_HOST_DEVICE (Int32 input_index,Int32 output_index)
516 * {
517 * output[output_index] = input[input_index];
518 * };
519 * Arcane::Accelerator::RunQueue* queue = ...;
520 * Arcane::Accelerator::GenericFilterer filterer(queue);
521 * filterer.applyWithIndex(input.size(), select_lambda, setter_lambda);
522 * Int32 nb_out = filterer.nbOutputElement();
523 * \endcode
524 */
525 template <typename SelectLambda, typename SetterLambda>
526 void applyWithIndex(Int32 nb_value, const SelectLambda& select_lambda,
527 const SetterLambda& setter_lambda, const TraceInfo& trace_info = TraceInfo())
528 {
529 _setCalled();
530 if (_checkEmpty(nb_value))
531 return;
532 impl::GenericFilteringBase* base_ptr = this;
534 impl::IndexIterator input_iter;
536 gf.apply(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
537 }
538
539 //! Nombre d'éléments en sortie.
541 {
542 m_is_already_called = false;
543 return _nbOutputElement();
544 }
545
546 private:
547
548 bool m_is_already_called = false;
549
550 private:
551
552 void _setCalled()
553 {
554 if (m_is_already_called)
555 ARCANE_FATAL("apply() has already been called for this instance");
556 m_is_already_called = true;
557 }
558 bool _checkEmpty(Int32 nb_value)
559 {
560 if (nb_value == 0) {
562 return true;
563 }
564 return false;
565 }
566};
567
568/*---------------------------------------------------------------------------*/
569/*---------------------------------------------------------------------------*/
570
571} // namespace Arcane::Accelerator
572
573/*---------------------------------------------------------------------------*/
574/*---------------------------------------------------------------------------*/
575
576#endif
577
578/*---------------------------------------------------------------------------*/
579/*---------------------------------------------------------------------------*/
#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.
Definition Filter.h:354
void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique un filtre.
Definition Filter.h:483
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.
Definition Filter.h:526
void apply(SmallSpan< const InputDataType > input, SmallSpan< OutputDataType > output, SmallSpan< const FlagType > flag)
Applique un filtre.
Definition Filter.h:408
void applyIf(SmallSpan< const DataType > input, SmallSpan< DataType > output, const SelectLambda &select_lambda, const TraceInfo &trace_info=TraceInfo())
Applique un filtre.
Definition Filter.h:459
Int32 nbOutputElement()
Nombre d'éléments en sortie.
Definition Filter.h:540
GenericFilterer(const RunQueue &queue)
Créé une instance.
Definition Filter.h:376
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:158
eExecutionPolicy executionPolicy() const
Politique d'exécution de la file.
Definition RunQueue.cc:168
DeviceStorage< int > m_device_nb_out_storage
Mémoire sur le device du nombre de valeurs filtrées.
Definition Filter.h:74
NumArray< Int32, MDDim1 > m_host_nb_out_storage
Mémoire hôte pour le nombre de valeurs filtrées.
Definition Filter.h:76
RunQueue m_queue
File d'exécution. Ne doit pas être nulle.
Definition Filter.h:70
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.
Definition Span.h:670
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:209
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:419
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.