Merge pull request #1207 from arcaneframework/dev/gg-add-internal-barrier-no-exception

Add method to synchronize stream without launching exception in case of error

Author: grospelliergilles

arcane/src/arcane/accelerator/core/IRunQueueStream.h

@@ -1,11 +1,11 @@
 // -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
 //-----------------------------------------------------------------------------
-// Copyright 2000-2022 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
+// Copyright 2000-2024 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: Apache-2.0
 //-----------------------------------------------------------------------------
 /*---------------------------------------------------------------------------*/
-/* IRunQueueStream.h                                           (C) 2000-2022 */
+/* IRunQueueStream.h                                           (C) 2000-2024 */
 /*                                                                           */
 /* Interface d'un flux d'exécution pour une RunQueue.                        */
 /*---------------------------------------------------------------------------*/
@@ -65,6 +65,9 @@ class ARCANE_ACCELERATOR_CORE_EXPORT IRunQueueStream
 
   //! Pointeur sur la structure interne dépendante de l'implémentation
   virtual void* _internalImpl() = 0;
+
+  //! Barrière sans exception. Retourne \a true en cas d'erreur
+  virtual bool _barrierNoException() = 0;
 };
 
 /*---------------------------------------------------------------------------*/

arcane/src/arcane/accelerator/core/RunQueueImpl.cc

@@ -1,11 +1,11 @@
 // -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
 //-----------------------------------------------------------------------------
-// Copyright 2000-2023 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
+// Copyright 2000-2024 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: Apache-2.0
 //-----------------------------------------------------------------------------
 /*---------------------------------------------------------------------------*/
-/* RunQueueImpl.cc                                             (C) 2000-2023 */
+/* RunQueueImpl.cc                                             (C) 2000-2024 */
 /*                                                                           */
 /* Gestion d'une file d'exécution sur accélérateur.                          */
 /*---------------------------------------------------------------------------*/
@@ -68,8 +68,13 @@ _release()
   // les commandes ne seront pas désallouées.
   // TODO: Regarder s'il ne faudrait pas plutôt indiquer cela à l'utilisateur
   // ou faire une erreur fatale.
-  if (!m_active_run_command_list.empty())
-    _internalBarrier();
+  if (!m_active_run_command_list.empty()){
+    if (!_internalStream()->_barrierNoException()){
+      _internalFreeRunningCommands();
+    }
+    else
+      std::cerr << "WARNING: Error in internal accelerator barrier\n";
+  }
   if (_isInPool())
     m_runner->_internalPutRunQueueImplInPool(this);
   else

arcane/src/arcane/accelerator/core/RunQueueRuntime.cc

@@ -1,11 +1,11 @@
 // -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
 //-----------------------------------------------------------------------------
-// Copyright 2000-2022 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
+// Copyright 2000-2024 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: Apache-2.0
 //-----------------------------------------------------------------------------
 /*---------------------------------------------------------------------------*/
-/* RunQueueRuntime.cc                                          (C) 2000-2022 */
+/* RunQueueRuntime.cc                                          (C) 2000-2024 */
 /*                                                                           */
 /* Implémentation d'un RunQueue pour une cible donnée.                       */
 /*---------------------------------------------------------------------------*/
@@ -53,6 +53,7 @@ class ARCANE_ACCELERATOR_CORE_EXPORT HostRunQueueStream
   }
   void prefetchMemory(const MemoryPrefetchArgs&) override {}
   void* _internalImpl() override { return nullptr; }
+  bool _barrierNoException() override { return false; }
 
  private:
 

arcane/src/arcane/accelerator/cuda/runtime/CudaAcceleratorRuntime.cc

@@ -148,6 +148,10 @@ class CudaRunQueueStream
     if (global_cupti_flush > 0)
       flushCupti();
   }
+  bool _barrierNoException() override
+  {
+    return (cudaStreamSynchronize(m_cuda_stream)!=CUDA_SUCCESS);
+  }
   void copyMemory(const MemoryCopyArgs& args) override
   {
     auto source_bytes = args.source().bytes();

arcane/src/arcane/accelerator/hip/runtime/HipAcceleratorRuntime.cc

@@ -90,6 +90,10 @@ class HipRunQueueStream
   {
     ARCANE_CHECK_HIP(hipStreamSynchronize(m_hip_stream));
   }
+  bool _barrierNoException() override
+  {
+    return hipStreamSynchronize(m_hip_stream) != HIP_SUCCESS;
+  }
   void copyMemory(const MemoryCopyArgs& args) override
   {
     auto r = hipMemcpyAsync(args.destination().data(), args.source().data(),

arcane/src/arcane/core/materials/ComponentItemInternal.h

@@ -437,7 +437,7 @@ _superItemBase(ConstituentItemIndex id) const
   return m_super_component_item_shared_info->_item(super_local_id);
 }
 
-inline matimpl::ConstituentItemBase ComponentItemSharedInfo::
+inline ARCCORE_HOST_DEVICE matimpl::ConstituentItemBase ComponentItemSharedInfo::
 _subItemBase(ConstituentItemIndex id,Int32 sub_index) const
 {
   ARCCORE_CHECK_RANGE(id.localId(), -1, m_storage_size);