Arcane  v3.16.0.0
Documentation développeur
Chargement...
Recherche...
Aucune correspondance
SyclAcceleratorRuntime.cc
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2025 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/* SyclAcceleratorRuntime.cc (C) 2000-2025 */
9/* */
10/* Runtime pour 'SYCL'. */
11/*---------------------------------------------------------------------------*/
12/*---------------------------------------------------------------------------*/
13
14#include "arcane/accelerator/sycl/SyclAccelerator.h"
15#include "arcane/accelerator/sycl/internal/SyclAcceleratorInternal.h"
16
17#include "arcane/utils/PlatformUtils.h"
18#include "arcane/utils/NotSupportedException.h"
19#include "arcane/utils/NotImplementedException.h"
20#include "arcane/utils/FatalErrorException.h"
21#include "arcane/utils/IMemoryRessourceMng.h"
22#include "arcane/utils/internal/IMemoryRessourceMngInternal.h"
23
24#include "arcane/accelerator/core/RunQueueBuildInfo.h"
25#include "arcane/accelerator/core/Memory.h"
26#include "arcane/accelerator/core/DeviceInfoList.h"
27#include "arcane/accelerator/core/RunQueue.h"
28#include "arcane/accelerator/core/DeviceMemoryInfo.h"
29#include "arcane/accelerator/core/NativeStream.h"
30#include "arcane/accelerator/core/internal/IRunnerRuntime.h"
31#include "arcane/accelerator/core/internal/RegisterRuntimeInfo.h"
32#include "arcane/accelerator/core/internal/IRunQueueStream.h"
33#include "arcane/accelerator/core/internal/IRunQueueEventImpl.h"
34
35#include <iostream>
36
37namespace Arcane::Accelerator::Sycl
38{
39
40using namespace Arccore;
41
42#define ARCANE_SYCL_FUNC_NOT_HANDLED \
43 std::cout << "WARNING: SYCL: function not handled " << A_FUNCINFO << "\n"
44
46
47/*---------------------------------------------------------------------------*/
48/*---------------------------------------------------------------------------*/
49
50class SyclRunQueueStream
52{
53 public:
54
55 SyclRunQueueStream(SyclRunnerRuntime* runtime, const RunQueueBuildInfo& bi);
56 ~SyclRunQueueStream() override
57 {
58 }
59
60 public:
61
62 void notifyBeginLaunchKernel([[maybe_unused]] impl::RunCommandImpl& c) override
63 {
64 return m_runtime->notifyBeginLaunchKernel();
65 }
67 {
68 return m_runtime->notifyEndLaunchKernel();
69 }
70 void barrier() override
71 {
72 m_sycl_stream->wait_and_throw();
73 }
74 bool _barrierNoException() override
75 {
76 m_sycl_stream->wait();
77 return false;
78 }
79 void copyMemory(const MemoryCopyArgs& args) override
80 {
81 auto source_bytes = args.source().bytes();
82 m_sycl_stream->memcpy(args.destination().data(), source_bytes.data(),
83 source_bytes.size());
84 if (!args.isAsync())
85 this->barrier();
86 }
87 void prefetchMemory([[maybe_unused]] const MemoryPrefetchArgs& args) override
88 {
89 auto source_bytes = args.source().bytes();
90 Int64 nb_byte = source_bytes.size();
91 if (nb_byte == 0)
92 return;
93 m_sycl_stream->prefetch(source_bytes.data(), nb_byte);
94 if (!args.isAsync())
95 this->barrier();
96 }
98 {
99 return impl::NativeStream(m_sycl_stream.get());
100 }
101
102 void _setSyclLastCommandEvent([[maybe_unused]] void* sycl_event_ptr) override
103 {
104 sycl::event last_event;
105 if (sycl_event_ptr)
106 last_event = *(reinterpret_cast<sycl::event*>(sycl_event_ptr));
107 m_last_command_event = last_event;
108 }
109
110 public:
111
112 static sycl::async_handler _getAsyncHandler()
113 {
114 auto err_handler = [](const sycl::exception_list& exceptions) {
115 std::ostringstream ostr;
116 ostr << "Error in SYCL runtime\n";
117 for (const std::exception_ptr& e : exceptions) {
118 try {
119 std::rethrow_exception(e);
120 }
121 catch (const sycl::exception& e) {
122 ostr << "SYCL exception: " << e.what() << "\n";
123 }
124 }
125 ARCANE_FATAL(ostr.str());
126 };
127 return err_handler;
128 }
129
131 sycl::event lastCommandEvent() { return m_last_command_event; }
132
133 public:
134
135 sycl::queue& trueStream() const
136 {
137 return *m_sycl_stream;
138 }
139
140 private:
141
142 impl::IRunnerRuntime* m_runtime;
143 std::unique_ptr<sycl::queue> m_sycl_stream;
144 sycl::event m_last_command_event;
145};
146
147/*---------------------------------------------------------------------------*/
148/*---------------------------------------------------------------------------*/
149
150class SyclRunQueueEvent
152{
153 public:
154
155 explicit SyclRunQueueEvent([[maybe_unused]] bool has_timer)
156 {
157 }
158 ~SyclRunQueueEvent() override
159 {
160 }
161
162 public:
163
164 // Enregistre l'événement au sein d'une RunQueue
165 void recordQueue([[maybe_unused]] impl::IRunQueueStream* stream) final
166 {
167 ARCANE_CHECK_POINTER(stream);
168 auto* rq = static_cast<SyclRunQueueStream*>(stream);
169 m_sycl_event = rq->lastCommandEvent();
170#if defined(__ADAPTIVECPP__)
171 m_recorded_stream = stream;
172 // TODO: Vérifier s'il faut faire quelque chose
173#elif defined(__INTEL_LLVM_COMPILER)
174 //m_sycl_event = rq->trueStream().ext_oneapi_submit_barrier();
175#else
176 ARCANE_THROW(NotSupportedException, "Only supported for AdaptiveCpp and Intel DPC++ implementation");
177#endif
178 }
179
180 void wait() final
181 {
182 //ARCANE_SYCL_FUNC_NOT_HANDLED;
183 // TODO: Vérifier ce que cela signifie exactement
184 m_sycl_event.wait();
185 }
186
187 void waitForEvent([[maybe_unused]] impl::IRunQueueStream* stream) final
188 {
189#if defined(__ADAPTIVECPP__)
190 auto* rq = static_cast<SyclRunQueueStream*>(stream);
191 m_sycl_event.wait(rq->trueStream().get_wait_list());
192#elif defined(__INTEL_LLVM_COMPILER)
193 std::vector<sycl::event> events;
194 events.push_back(m_sycl_event);
195 auto* rq = static_cast<SyclRunQueueStream*>(stream);
196 rq->trueStream().ext_oneapi_submit_barrier(events);
197#else
198 ARCANE_THROW(NotSupportedException, "Only supported for AdaptiveCpp and Intel DPC++ implementation");
199#endif
200 }
201
202 Int64 elapsedTime([[maybe_unused]] IRunQueueEventImpl* start_event) final
203 {
204 ARCANE_CHECK_POINTER(start_event);
205 // Il faut prendre l'évènement de début car on est certain qu'il contient
206 // la bonne valeur de 'sycl::event'.
207 sycl::event event = (static_cast<SyclRunQueueEvent*>(start_event))->m_sycl_event;
208 // Si pas d'évènement associé, on ne fait rien pour éviter une exception
209 if (event==sycl::event())
210 return 0;
211
212 bool is_submitted = event.get_info<sycl::info::event::command_execution_status>() == sycl::info::event_command_status::complete;
213 if (!is_submitted)
214 return 0;
215 Int64 start = event.get_profiling_info<sycl::info::event_profiling::command_start>();
216 Int64 end = event.get_profiling_info<sycl::info::event_profiling::command_end>();
217 return (end - start);
218 }
219
220 bool hasPendingWork() final
221 {
222 ARCANE_THROW(NotImplementedException,"hasPendingWork()");
223 }
224
225 private:
226
227 sycl::event m_sycl_event;
228 impl::IRunQueueStream* m_recorded_stream = nullptr;
229};
230
231/*---------------------------------------------------------------------------*/
232/*---------------------------------------------------------------------------*/
233
236{
237 friend class SyclRunQueueStream;
238
239 public:
240
241 void notifyBeginLaunchKernel() override
242 {
243 }
244 void notifyEndLaunchKernel() override
245 {
246 }
247 void barrier() override
248 {
249 // TODO Faire le wait sur la file par défaut n'est pas strictement équivalent
250 // à la barrière en CUDA qui synchronize tout le device.
251 m_default_queue->wait();
252 }
253 eExecutionPolicy executionPolicy() const override
254 {
256 }
257 impl::IRunQueueStream* createStream(const RunQueueBuildInfo& bi) override
258 {
259 return new SyclRunQueueStream(this, bi);
260 }
261 impl::IRunQueueEventImpl* createEventImpl() override
262 {
263 return new SyclRunQueueEvent(false);
264 }
265 impl::IRunQueueEventImpl* createEventImplWithTimer() override
266 {
267 return new SyclRunQueueEvent(true);
268 }
269 void setMemoryAdvice([[maybe_unused]] ConstMemoryView buffer, [[maybe_unused]] eMemoryAdvice advice,
270 [[maybe_unused]] DeviceId device_id) override
271 {
272 }
273 void unsetMemoryAdvice([[maybe_unused]] ConstMemoryView buffer,
274 [[maybe_unused]] eMemoryAdvice advice, [[maybe_unused]] DeviceId device_id) override
275 {
276 }
277
278 void setCurrentDevice([[maybe_unused]] DeviceId device_id) final
279 {
280 ARCANE_SYCL_FUNC_NOT_HANDLED;
281 }
282 const IDeviceInfoList* deviceInfoList() override { return &m_device_info_list; }
283
284 void getPointerAttribute(PointerAttribute& attribute, const void* ptr) override
285 {
286 sycl::usm::alloc sycl_mem_type = sycl::get_pointer_type(ptr, *m_default_context);
287 ePointerMemoryType mem_type = ePointerMemoryType::Unregistered;
288 const void* host_ptr = nullptr;
289 const void* device_ptr = nullptr;
290 if (sycl_mem_type == sycl::usm::alloc::host) {
291 // HostPinned. Doit être accessible depuis le device mais
292 //
293 mem_type = ePointerMemoryType::Host;
294 host_ptr = ptr;
295 // TODO: Regarder comment récupérer la valeur
296 device_ptr = ptr;
297 }
298 else if (sycl_mem_type == sycl::usm::alloc::device) {
299 mem_type = ePointerMemoryType::Device;
300 device_ptr = ptr;
301 }
302 else if (sycl_mem_type == sycl::usm::alloc::shared) {
303 mem_type = ePointerMemoryType::Managed;
304 // TODO: pour l'instant on remplit avec le pointeur car on ne sait
305 // pas comment récupérer l'info.
306 host_ptr = ptr;
307 device_ptr = ptr;
308 }
309 // TODO: à corriger
310 Int32 device_id = 0;
311 _fillPointerAttribute(attribute, mem_type, device_id, ptr, device_ptr, host_ptr);
312 }
313
314 DeviceMemoryInfo getDeviceMemoryInfo(DeviceId device_id) override
315 {
316 return {};
317 }
318
319 void fillDevicesAndSetDefaultQueue(bool is_verbose);
320 sycl::queue& defaultQueue() const { return *m_default_queue; }
321 sycl::device& defaultDevice() const { return *m_default_device; }
322
323 private:
324
325 impl::DeviceInfoList m_device_info_list;
326 std::unique_ptr<sycl::device> m_default_device;
327 std::unique_ptr<sycl::context> m_default_context;
328 std::unique_ptr<sycl::queue> m_default_queue;
329
330 private:
331
332 void _init(sycl::device& device)
333 {
334 m_default_device = std::make_unique<sycl::device>(device);
335 m_default_queue = std::make_unique<sycl::queue>(device);
336 m_default_context = std::make_unique<sycl::context>(device);
337 }
338};
339
340/*---------------------------------------------------------------------------*/
341/*---------------------------------------------------------------------------*/
342
343SyclRunQueueStream::
344SyclRunQueueStream(SyclRunnerRuntime* runtime, const RunQueueBuildInfo& bi)
345: m_runtime(runtime)
346{
347 sycl::device& d = runtime->defaultDevice();
348 // Indique que les commandes lancées sont implicitement exécutées les
349 // unes derrière les autres.
350 auto queue_property = sycl::property::queue::in_order();
351 // Pour le profiling
352 auto profiling_property = sycl::property::queue::enable_profiling();
353 sycl::property_list queue_properties(queue_property, profiling_property);
354
355 // Gestionnaire d'erreur.
356 sycl::async_handler err_handler;
357 err_handler = _getAsyncHandler();
358 if (bi.isDefault())
359 m_sycl_stream = std::make_unique<sycl::queue>(d, err_handler, queue_properties);
360 else {
361 ARCANE_SYCL_FUNC_NOT_HANDLED;
362 m_sycl_stream = std::make_unique<sycl::queue>(d, err_handler, queue_properties);
363 }
364}
365
366/*---------------------------------------------------------------------------*/
367/*---------------------------------------------------------------------------*/
368
369void SyclRunnerRuntime::
370fillDevicesAndSetDefaultQueue(bool is_verbose)
371{
372 if (is_verbose){
373 for (auto platform : sycl::platform::get_platforms()) {
374 std::cout << "Platform: "
375 << platform.get_info<sycl::info::platform::name>()
376 << std::endl;
377 }
378 }
379
380 sycl::device device{ sycl::gpu_selector_v };
381 if (is_verbose)
382 std::cout << "\nDevice: " << device.get_info<sycl::info::device::name>()
383 << "\nVersion=" << device.get_info<sycl::info::device::version>()
384 << std::endl;
385 // Pour l'instant, on prend comme file par défaut la première trouvée
386 // et on ne considère qu'un seul device accessible.
387 _init(device);
388
389 DeviceInfo device_info;
390 device_info.setDescription("No description info");
391 device_info.setDeviceId(DeviceId(0));
392 device_info.setName(device.get_info<sycl::info::device::name>());
393 m_device_info_list.addDevice(device_info);
394}
395
396/*---------------------------------------------------------------------------*/
397/*---------------------------------------------------------------------------*/
398
400: public IMemoryCopier
401{
402 void copy(ConstMemoryView from, eMemoryRessource from_mem,
403 MutableMemoryView to, eMemoryRessource to_mem,
404 const RunQueue* queue) override;
405};
406
407/*---------------------------------------------------------------------------*/
408/*---------------------------------------------------------------------------*/
409
410} // namespace Arcane::Accelerator::Sycl
411
412namespace
413{
415Arcane::Accelerator::Sycl::SyclMemoryCopier global_sycl_memory_copier;
416} // namespace
417
418/*---------------------------------------------------------------------------*/
419/*---------------------------------------------------------------------------*/
420
421namespace Arcane::Accelerator::Sycl
422{
423
424/*---------------------------------------------------------------------------*/
425/*---------------------------------------------------------------------------*/
426
427void SyclMemoryCopier::
428copy(ConstMemoryView from, [[maybe_unused]] eMemoryRessource from_mem,
429 MutableMemoryView to, [[maybe_unused]] eMemoryRessource to_mem,
430 const RunQueue* queue)
431{
432 if (queue) {
433 queue->copyMemory(MemoryCopyArgs(to.bytes(), from.bytes()).addAsync(queue->isAsync()));
434 return;
435 }
436 sycl::queue& q = global_sycl_runtime.defaultQueue();
437 q.memcpy(to.data(), from.data(), from.bytes().size()).wait();
438}
439
440} // namespace Arcane::Accelerator::Sycl
441
442/*---------------------------------------------------------------------------*/
443/*---------------------------------------------------------------------------*/
444
445// Cette fonction est le point d'entrée utilisé lors du chargement
446// dynamique de cette bibliothèque
447extern "C" ARCANE_EXPORT void
448arcaneRegisterAcceleratorRuntimesycl(Arcane::Accelerator::RegisterRuntimeInfo& init_info)
449{
450 using namespace Arcane;
451 using namespace Arcane::Accelerator::Sycl;
452 Arcane::Accelerator::impl::setUsingSYCLRuntime(true);
453 Arcane::Accelerator::impl::setSYCLRunQueueRuntime(&global_sycl_runtime);
456 mrm->setIsAccelerator(true);
457 mrm->setAllocator(eMemoryRessource::UnifiedMemory, getSyclUnifiedMemoryAllocator());
458 mrm->setAllocator(eMemoryRessource::HostPinned, getSyclHostPinnedMemoryAllocator());
459 mrm->setAllocator(eMemoryRessource::Device, getSyclDeviceMemoryAllocator());
460 mrm->setCopier(&global_sycl_memory_copier);
461 global_sycl_runtime.fillDevicesAndSetDefaultQueue(init_info.isVerbose());
462 setSyclMemoryQueue(global_sycl_runtime.defaultQueue());
463}
464
465/*---------------------------------------------------------------------------*/
466/*---------------------------------------------------------------------------*/
#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_THROW(exception_class,...)
Macro pour envoyer une exception avec formattage.
#define ARCANE_FATAL(...)
Macro envoyant une exception FatalErrorException.
Identifiant d'un composant du système.
Definition DeviceId.h:33
Information mémoire d'un accélérateur.
Interface d'une liste de devices.
Arguments pour la copie mémoire.
Definition Memory.h:63
Arguments pour le préfetching mémoire.
Definition Memory.h:125
Informations sur une adresse mémoire.
Informations pour initialiser le runtime accélérateur.
Informations pour créer une RunQueue.
bool isDefault() const
Indique si l'instance a uniquement les valeurs par défaut.
File d'exécution pour un accélérateur.
bool isAsync() const
Indique si la file d'exécution est asynchrone.
Definition RunQueue.cc:320
void copyMemory(const MemoryCopyArgs &args) const
Copie des informations entre deux zones mémoires.
Definition RunQueue.cc:237
void copy(ConstMemoryView from, eMemoryRessource from_mem, MutableMemoryView to, eMemoryRessource to_mem, const RunQueue *queue) override
Copie les données de from vers to avec la queue queue.
impl::NativeStream nativeStream() override
Pointeur sur la structure interne dépendante de l'implémentation.
void barrier() override
Bloque jusqu'à ce que toutes les actions associées à cette file soient terminées.
sycl::event lastCommandEvent()
Évènement correspondant à la dernière commande.
void prefetchMemory(const MemoryPrefetchArgs &args) override
Effectue un pré-chargement d'une zone mémoire.
void copyMemory(const MemoryCopyArgs &args) override
Effectue une copie entre deux zones mémoire.
bool _barrierNoException() override
Barrière sans exception. Retourne true en cas d'erreur.
void _setSyclLastCommandEvent(void *sycl_event_ptr) override
Pour SYCL, positionne l'évènement associé à la dernière commande exécutée.
void notifyEndLaunchKernel(impl::RunCommandImpl &) override
Notification de fin de lancement de la commande.
void notifyBeginLaunchKernel(impl::RunCommandImpl &c) override
Notification avant le lancement de la commande.
Interface d'une liste de devices.
Interface de l'implémentation d'un évènement.
Interface d'un flux d'exécution pour une RunQueue.
Interface du runtime associé à un accélérateur.
Type opaque pour encapsuler une 'stream' native.
Implémentation d'une commande pour accélérateur.
Interface pour les copies mémoire avec support des accélérateurs.
virtual IMemoryRessourceMngInternal * _internal()=0
Interface interne.
Partie interne à Arcane de 'IMemoryRessourceMng'.
virtual void setAllocator(eMemoryRessource r, IMemoryAllocator *allocator)=0
Positionne l'allocateur pour la ressource r.
virtual void setCopier(IMemoryCopier *copier)=0
Positionne l'instance gérant les copies.
virtual void setIsAccelerator(bool v)=0
Indique si un accélérateur est disponible.
Exception lorsqu'une fonction n'est pas implémentée.
Exception lorsqu'une opération n'est pas supportée.
eMemoryAdvice
Conseils pour la gestion mémoire.
Definition Memory.h:36
ePointerMemoryType
Type de mémoire pour un pointeur.
eExecutionPolicy
Politique d'exécution pour un Runner.
@ SYCL
Politique d'exécution utilisant l'environnement SYCL.
Espace de nom pour les fonctions dépendant de la plateforme.
IMemoryRessourceMng * getDataMemoryRessourceMng()
Gestionnaire de ressource mémoire pour les données.
IMemoryAllocator * setAcceleratorHostMemoryAllocator(IMemoryAllocator *a)
Positionne l'allocateur spécifique pour les accélérateurs.
-*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
@ HostPinned
Alloue sur l'hôte.
@ UnifiedMemory
Alloue en utilisant la mémoire unifiée.
@ Device
Alloue sur le device.
Arcane::eMemoryResource eMemoryRessource
Typedef pour la version Arcane historique (avec 2's')