Merge pull request #1882 from arcaneframework/dev/gg-add-in-place-filtering

Add specific implementation of 'GenericFiltering' when input and output are the same

Author: grospelliergilles

arcane/src/arcane/accelerator/GenericFilterer.h

@@ -178,7 +178,8 @@ class GenericFilteringFlag
 {
  public:
 
-  void apply(GenericFilteringBase& s, SmallSpan<const DataType> input, SmallSpan<OutputDataType> output, SmallSpan<const FlagType> flag)
+  void apply(GenericFilteringBase& s, SmallSpan<const DataType> input,
+             SmallSpan<OutputDataType> output, SmallSpan<const FlagType> flag)
   {
     const Int32 nb_item = input.size();
     if (output.size() != nb_item)
@@ -264,7 +265,13 @@ class GenericFilteringIf
 {
  public:
 
-  template <typename SelectLambda, typename InputIterator, typename OutputIterator>
+  /*!
+   * \brief Applique le filtre.
+   *
+   * Si \a InPlace est vrai, alors OutputIterator vaut InputIterator et on
+   * met à jour directement \a input_iter.
+   */
+  template <bool InPlace, typename SelectLambda, typename InputIterator, typename OutputIterator>
   void apply(GenericFilteringBase& s, Int32 nb_item, InputIterator input_iter, OutputIterator output_iter,
              const SelectLambda& select_lambda, const TraceInfo& trace_info)
   {
@@ -282,15 +289,26 @@ class GenericFilteringIf
       cudaStream_t stream = impl::CudaUtils::toNativeStream(queue);
       // Premier appel pour connaitre la taille pour l'allocation
       int* nb_out_ptr = nullptr;
-      ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(nullptr, temp_storage_size,
-                                                input_iter, output_iter, nb_out_ptr, nb_item,
-                                                select_lambda, stream));
+      if constexpr (InPlace)
+        ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(nullptr, temp_storage_size,
+                                                  input_iter, nb_out_ptr, nb_item,
+                                                  select_lambda, stream));
+      else
+        ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(nullptr, temp_storage_size,
+                                                  input_iter, output_iter, nb_out_ptr, nb_item,
+                                                  select_lambda, stream));
 
       s._allocateTemporaryStorage(temp_storage_size);
       nb_out_ptr = s._getDeviceNbOutPointer();
-      ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(s.m_algo_storage.address(), temp_storage_size,
-                                                input_iter, output_iter, nb_out_ptr, nb_item,
-                                                select_lambda, stream));
+      if constexpr (InPlace)
+        ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(s.m_algo_storage.address(), temp_storage_size,
+                                                  input_iter, nb_out_ptr, nb_item,
+                                                  select_lambda, stream));
+      else
+        ARCANE_CHECK_CUDA(::cub::DeviceSelect::If(s.m_algo_storage.address(), temp_storage_size,
+                                                  input_iter, output_iter, nb_out_ptr, nb_item,
+                                                  select_lambda, stream));
+
       s._copyDeviceNbOutToHostNbOut();
     } break;
 #endif
@@ -300,9 +318,11 @@ class GenericFilteringIf
       // Premier appel pour connaitre la taille pour l'allocation
       hipStream_t stream = impl::HipUtils::toNativeStream(queue);
       int* nb_out_ptr = nullptr;
+      // NOTE: il n'y a pas de version spécifique de 'select' en-place.
+      // A priori il est possible que \a input_iter et \a output_iter
+      // aient la même valeur.
       ARCANE_CHECK_HIP(rocprim::select(nullptr, temp_storage_size, input_iter, output_iter,
                                        nb_out_ptr, nb_item, select_lambda, stream));
-
       s._allocateTemporaryStorage(temp_storage_size);
       nb_out_ptr = s._getDeviceNbOutPointer();
       ARCANE_CHECK_HIP(rocprim::select(s.m_algo_storage.address(), temp_storage_size, input_iter, output_iter,
@@ -318,7 +338,7 @@ class GenericFilteringIf
     case eExecutionPolicy::Thread:
       if (nb_item > 500) {
         MultiThreadAlgo scanner;
-        Int32 v = scanner.doFilter(launch_info.loopRunInfo(), nb_item, input_iter, output_iter, select_lambda);
+        Int32 v = scanner.doFilter<InPlace>(launch_info.loopRunInfo(), nb_item, input_iter, output_iter, select_lambda);
         s.m_host_nb_out_storage[0] = v;
         break;
       }
@@ -433,6 +453,7 @@ class GenericFilterer
    * Filtre tous les éléments de \a input pour lesquels \a select_lambda vaut \a true et
    * remplit \a output avec les valeurs filtrées. \a output doit avoir une taille assez
    * grande pour contenir tous les éléments filtrés.
+   * Les zones mémoire associées à \a input et \a output ne doivent pas se chevaucher.
    *
    * \a select_lambda doit avoir un opérateur `ARCCORE_HOST_DEVICE bool operator()(const DataType& v) const`.
    *
@@ -468,13 +489,36 @@ class GenericFilterer
     const Int32 nb_value = input.size();
     if (output.size() != nb_value)
       ARCANE_FATAL("Sizes are not equals: input={0} output={1}", nb_value, output.size());
+    if (input.data() == output.data())
+      ARCANE_FATAL("Input and Output are the same. Use in place overload instead");
+    _setCalled();
+    if (_checkEmpty(nb_value))
+      return;
+    impl::GenericFilteringBase* base_ptr = this;
+    impl::GenericFilteringIf gf;
+    gf.apply<false>(*base_ptr, nb_value, input.data(), output.data(), select_lambda, trace_info);
+  }
 
+  /*!
+   * \brief Applique un filtre en place.
+   *
+   * Cette méthode est identique à applyIf(SmallSpan<const DataType>, SmallSpan<DataType>,
+   * const SelectLambda&, const TraceInfo& trace_info) mais les valeurs filtrées sont
+   * directement recopié dans le tableau \a input_output.
+   */
+  template <typename DataType, typename SelectLambda>
+  void applyIf(SmallSpan<DataType> input_output, const SelectLambda& select_lambda,
+               const TraceInfo& trace_info = TraceInfo())
+  {
+    const Int32 nb_value = input_output.size();
+    if (nb_value <= 0)
+      return;
     _setCalled();
     if (_checkEmpty(nb_value))
       return;
     impl::GenericFilteringBase* base_ptr = this;
     impl::GenericFilteringIf gf;
-    gf.apply(*base_ptr, nb_value, input.data(), output.data(), select_lambda, trace_info);
+    gf.apply<true>(*base_ptr, nb_value, input_output.data(), input_output.data(), select_lambda, trace_info);
   }
 
   /*!
@@ -484,6 +528,8 @@ class GenericFilterer
    * SmallSpan<DataType> output, const SelectLambda& select_lambda) mais permet de spécifier un
    * itérateur \a input_iter pour l'entrée et \a output_iter pour la sortie.
    * Le nombre d'entité en entrée est donné par \a nb_value.
+   *
+   * Les zones mémoire associées à \a input_iter et \a output_iter ne doivent pas se chevaucher.
    */
   template <typename InputIterator, typename OutputIterator, typename SelectLambda>
   void applyIf(Int32 nb_value, InputIterator input_iter, OutputIterator output_iter,
@@ -494,7 +540,7 @@ class GenericFilterer
       return;
     impl::GenericFilteringBase* base_ptr = this;
     impl::GenericFilteringIf gf;
-    gf.apply(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
+    gf.apply<false>(*base_ptr, nb_value, input_iter, output_iter, select_lambda, trace_info);
   }
 
   /*!
@@ -527,6 +573,8 @@ class GenericFilterer
    * filterer.applyWithIndex(input.size(), select_lambda, setter_lambda);
    * Int32 nb_out = filterer.nbOutputElement();
    * \endcode
+   *
+   * Les zones mémoire associées aux valeurs d'entrée et de sortie ne doivent pas se chevaucher.
    */
   template <typename SelectLambda, typename SetterLambda>
   void applyWithIndex(Int32 nb_value, const SelectLambda& select_lambda,
@@ -539,7 +587,7 @@ class GenericFilterer
     impl::GenericFilteringIf gf;
     impl::IndexIterator input_iter;
     impl::SetterLambdaIterator<SetterLambda> out(setter_lambda);
-    gf.apply(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
+    gf.apply<false>(*base_ptr, nb_value, input_iter, out, select_lambda, trace_info);
   }
 
   //! Nombre d'éléments en sortie.

arcane/src/arcane/accelerator/MultiThreadAlgo.h

@@ -130,7 +130,7 @@ class MultiThreadAlgo
     Arcane::arcaneParallelFor(0, nb_block, run_info, final_sum_func);
   }
 
-  template <typename InputIterator, typename OutputIterator, typename SelectLambda>
+  template <bool InPlace, typename InputIterator, typename OutputIterator, typename SelectLambda>
   Int32 doFilter(ForLoopRunInfo run_info, Int32 nb_value,
                  InputIterator input, OutputIterator output,
                  SelectLambda select_lambda)
@@ -215,10 +215,9 @@ class MultiThreadAlgo
       }
     };
 
-    // Pour l'instant il est possible que l'entrée et la sortie
-    // se chevauchent. Dans ce cas on fait le remplissage en séquentiel.
-    const bool may_input_and_output_overlap = true;
-    if (may_input_and_output_overlap)
+    // Si l'entrée et la sortie sont les mêmes, on fait le remplissage en séquentiel.
+    // TODO: faire en parallèle.
+    if (InPlace)
       filter_func(0, nb_block);
     else
       Arcane::arcaneParallelFor(0, nb_block, run_info, filter_func);

arcane/src/arcane/materials/IncrementalComponentModifier_Accelerator.cc

@@ -362,14 +362,13 @@ _removeItemsInGroup(ItemGroup cells, SmallSpan<const Int32> removed_ids)
 
     // Lors de l'application du filtre, le tableau d'entrée et de sortie
     // est le même (c'est normalement supporté par le GenericFilterer).
-    SmallSpan<const Int32> input_ids(items_local_id);
-    SmallSpan<Int32> output_ids_view(items_local_id);
+    SmallSpan<Int32> input_ids(items_local_id);
     SmallSpan<const bool> filtered_cells(m_work_info.removedCells());
     Accelerator::GenericFilterer filterer(m_queue);
     auto select_filter = [=] ARCCORE_HOST_DEVICE(Int32 local_id) -> bool {
       return !filtered_cells[local_id];
     };
-    filterer.applyIf(input_ids, output_ids_view, select_filter, A_FUNCINFO);
+    filterer.applyIf(input_ids, select_filter, A_FUNCINFO);
 
     Int32 current_nb_item = items_local_id.size();
     Int32 nb_remaining = filterer.nbOutputElement();