RunCommandLaunchInfo.h

// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
//-----------------------------------------------------------------------------
// 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
//-----------------------------------------------------------------------------
/*---------------------------------------------------------------------------*/
/* RunCommandLaunchInfo.h                                      (C) 2000-2024 */
/*                                                                           */
/* Informations pour l'exécution d'une 'RunCommand'.                         */
/*---------------------------------------------------------------------------*/
#ifndef ARCANE_ACCELERATOR_RUNCOMMANDLAUNCHINFO_H
#define ARCANE_ACCELERATOR_RUNCOMMANDLAUNCHINFO_H
/*---------------------------------------------------------------------------*/
/*---------------------------------------------------------------------------*/
 
#include "arcane/utils/CheckedConvert.h"
#include "arcane/utils/ConcurrencyUtils.h"
#include "arcane/utils/Profiling.h"
 
#include "arcane/accelerator/AcceleratorGlobal.h"
 
/*---------------------------------------------------------------------------*/
/*---------------------------------------------------------------------------*/
 
namespace Arcane::Accelerator::impl
{
 
/*---------------------------------------------------------------------------*/
/*---------------------------------------------------------------------------*/
class ARCANE_ACCELERATOR_EXPORT RunCommandLaunchInfo
{
  template <typename SyclKernel, typename Lambda, typename LambdaArgs, typename... ReducerArgs>
  friend void _applyKernelSYCL(impl::RunCommandLaunchInfo& launch_info, SyclKernel kernel, Lambda& func,
                               const LambdaArgs& args, const ReducerArgs&... reducer_args);
 
 public:
 
  struct ThreadBlockInfo
  {
    int nb_block_per_grid = 0;
    int nb_thread_per_block = 0;
  };
 
 public:
 
  RunCommandLaunchInfo(RunCommand& command, Int64 total_loop_size);
  ~RunCommandLaunchInfo();
  RunCommandLaunchInfo(const RunCommandLaunchInfo&) = delete;
  RunCommandLaunchInfo operator=(const RunCommandLaunchInfo&) = delete;
 
 public:
 
  eExecutionPolicy executionPolicy() const { return m_exec_policy; }
 
  void beginExecute();
 
  void endExecute();
 
  ThreadBlockInfo threadBlockInfo() const { return m_thread_block_info; }
 
  ParallelLoopOptions computeParallelLoopOptions() const;
 
  void computeLoopRunInfo();
 
  const ForLoopRunInfo& loopRunInfo() const { return m_loop_run_info; }
 
  Int64 totalLoopSize() const { return m_total_loop_size; }
 
 public:
 
  void* _internalStreamImpl();
 
 private:
 
  RunCommand& m_command;
  bool m_has_exec_begun = false;
  bool m_is_notify_end_kernel_done = false;
  IRunnerRuntime* m_runtime = nullptr;
  IRunQueueStream* m_queue_stream = nullptr;
  eExecutionPolicy m_exec_policy = eExecutionPolicy::Sequential;
  ThreadBlockInfo m_thread_block_info;
  ForLoopRunInfo m_loop_run_info;
  Int64 m_total_loop_size = 0;
 
 private:
 
  void _begin();
  void _doEndKernelLaunch();
  ThreadBlockInfo _computeThreadBlockInfo() const;
 
 private:
 
  // Pour SYCL: enregistre l'évènement associé à la dernière commande de la file
  // \a sycl_event_ptr est de type 'sycl::event*'.
  void _addSyclEvent(void* sycl_event_ptr);
};
 
/*---------------------------------------------------------------------------*/
/*---------------------------------------------------------------------------*/
 
} // End namespace Arcane::Accelerator::impl
 
/*---------------------------------------------------------------------------*/
/*---------------------------------------------------------------------------*/
 
#endif