Arcane  v4.1.0.0
Documentation développeur
Chargement...
Recherche...
Aucune correspondance
HipAcceleratorRuntime.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/* HipAcceleratorRuntime.cc (C) 2000-2025 */
9/* */
10/* Runtime pour 'HIP'. */
11/*---------------------------------------------------------------------------*/
12/*---------------------------------------------------------------------------*/
13
14#include "arcane/accelerator/hip/HipAccelerator.h"
15
16#include "arcane/utils/PlatformUtils.h"
17#include "arcane/utils/Array.h"
18#include "arcane/utils/TraceInfo.h"
19#include "arcane/utils/FatalErrorException.h"
20#include "arcane/utils/NotImplementedException.h"
21#include "arcane/utils/IMemoryRessourceMng.h"
22#include "arcane/utils/OStringStream.h"
23#include "arcane/utils/internal/MemoryUtilsInternal.h"
24#include "arcane/utils/internal/IMemoryRessourceMngInternal.h"
25
26#include "arcane/accelerator/core/RunQueueBuildInfo.h"
27#include "arcane/accelerator/core/Memory.h"
28#include "arcane/accelerator/core/DeviceInfoList.h"
29#include "arcane/accelerator/core/RunQueue.h"
30#include "arcane/accelerator/core/DeviceMemoryInfo.h"
31#include "arcane/accelerator/core/NativeStream.h"
32#include "arcane/accelerator/core/internal/IRunnerRuntime.h"
33#include "arcane/accelerator/core/internal/RegisterRuntimeInfo.h"
34#include "arcane/accelerator/core/internal/IRunQueueStream.h"
35#include "arcane/accelerator/core/internal/IRunQueueEventImpl.h"
36#include "arcane/accelerator/core/internal/RunCommandImpl.h"
37
38#include <iostream>
39
40#ifdef ARCANE_HAS_ROCTX
41#include <roctx.h>
42#endif
43
44using namespace Arccore;
45
46namespace Arcane::Accelerator::Hip
47{
48
49/*---------------------------------------------------------------------------*/
50/*---------------------------------------------------------------------------*/
51
52class HipRunQueueStream
54{
55 public:
56
57 HipRunQueueStream(impl::IRunnerRuntime* runtime, const RunQueueBuildInfo& bi)
58 : m_runtime(runtime)
59 {
60 if (bi.isDefault())
61 ARCANE_CHECK_HIP(hipStreamCreate(&m_hip_stream));
62 else {
63 int priority = bi.priority();
64 ARCANE_CHECK_HIP(hipStreamCreateWithPriority(&m_hip_stream, hipStreamDefault, priority));
65 }
66 }
67 ~HipRunQueueStream() override
68 {
69 ARCANE_CHECK_HIP_NOTHROW(hipStreamDestroy(m_hip_stream));
70 }
71
72 public:
73
74 void notifyBeginLaunchKernel([[maybe_unused]] impl::RunCommandImpl& c) override
75 {
76#ifdef ARCANE_HAS_ROCTX
77 auto kname = c.kernelName();
78 if (kname.empty())
79 roctxRangePush(c.traceInfo().name());
80 else
81 roctxRangePush(kname.localstr());
82#endif
83 return m_runtime->notifyBeginLaunchKernel();
84 }
86 {
87#ifdef ARCANE_HAS_ROCTX
88 roctxRangePop();
89#endif
90 return m_runtime->notifyEndLaunchKernel();
91 }
92 void barrier() override
93 {
94 ARCANE_CHECK_HIP(hipStreamSynchronize(m_hip_stream));
95 }
96 bool _barrierNoException() override
97 {
98 return hipStreamSynchronize(m_hip_stream) != hipSuccess;
99 }
100 void copyMemory(const MemoryCopyArgs& args) override
101 {
102 auto r = hipMemcpyAsync(args.destination().data(), args.source().data(),
103 args.source().bytes().size(), hipMemcpyDefault, m_hip_stream);
104 ARCANE_CHECK_HIP(r);
105 if (!args.isAsync())
106 barrier();
107 }
108 void prefetchMemory(const MemoryPrefetchArgs& args) override
109 {
110 auto src = args.source().bytes();
111 if (src.size()==0)
112 return;
113 DeviceId d = args.deviceId();
114 int device = hipCpuDeviceId;
115 if (!d.isHost())
116 device = d.asInt32();
117 auto r = hipMemPrefetchAsync(src.data(), src.size(), device, m_hip_stream);
118 ARCANE_CHECK_HIP(r);
119 if (!args.isAsync())
120 barrier();
121 }
123 {
124 return impl::NativeStream(&m_hip_stream);
125 }
126
127 public:
128
129 hipStream_t trueStream() const
130 {
131 return m_hip_stream;
132 }
133
134 private:
135
136 impl::IRunnerRuntime* m_runtime;
137 hipStream_t m_hip_stream;
138};
139
140/*---------------------------------------------------------------------------*/
141/*---------------------------------------------------------------------------*/
142
143class HipRunQueueEvent
145{
146 public:
147
148 explicit HipRunQueueEvent(bool has_timer)
149 {
150 if (has_timer)
151 ARCANE_CHECK_HIP(hipEventCreate(&m_hip_event));
152 else
153 ARCANE_CHECK_HIP(hipEventCreateWithFlags(&m_hip_event, hipEventDisableTiming));
154 }
155 ~HipRunQueueEvent() override
156 {
157 ARCANE_CHECK_HIP_NOTHROW(hipEventDestroy(m_hip_event));
158 }
159
160 public:
161
162 // Enregistre l'événement au sein d'une RunQueue
163 void recordQueue(impl::IRunQueueStream* stream) final
164 {
165 auto* rq = static_cast<HipRunQueueStream*>(stream);
166 ARCANE_CHECK_HIP(hipEventRecord(m_hip_event, rq->trueStream()));
167 }
168
169 void wait() final
170 {
171 ARCANE_CHECK_HIP(hipEventSynchronize(m_hip_event));
172 }
173
174 void waitForEvent(impl::IRunQueueStream* stream) final
175 {
176 auto* rq = static_cast<HipRunQueueStream*>(stream);
177 ARCANE_CHECK_HIP(hipStreamWaitEvent(rq->trueStream(), m_hip_event, 0));
178 }
179
180 Int64 elapsedTime(IRunQueueEventImpl* from_event) final
181 {
182 auto* true_from_event = static_cast<HipRunQueueEvent*>(from_event);
183 ARCANE_CHECK_POINTER(true_from_event);
184 float time_in_ms = 0.0;
185 ARCANE_CHECK_HIP(hipEventElapsedTime(&time_in_ms, true_from_event->m_hip_event, m_hip_event));
186 double x = time_in_ms * 1.0e6;
187 Int64 nano_time = static_cast<Int64>(x);
188 return nano_time;
189 }
190
191 bool hasPendingWork() final
192 {
193 hipError_t v = hipEventQuery(m_hip_event);
194 if (v == hipErrorNotReady)
195 return true;
196 ARCANE_CHECK_HIP(v);
197 return false;
198 }
199
200 private:
201
202 hipEvent_t m_hip_event;
203};
204
205/*---------------------------------------------------------------------------*/
206/*---------------------------------------------------------------------------*/
207
210{
211 public:
212
213 ~HipRunnerRuntime() override = default;
214
215 public:
216
217 void notifyBeginLaunchKernel() override
218 {
219 ++m_nb_kernel_launched;
220 if (m_is_verbose)
221 std::cout << "BEGIN HIP KERNEL!\n";
222 }
223 void notifyEndLaunchKernel() override
224 {
225 ARCANE_CHECK_HIP(hipGetLastError());
226 if (m_is_verbose)
227 std::cout << "END HIP KERNEL!\n";
228 }
229 void barrier() override
230 {
231 ARCANE_CHECK_HIP(hipDeviceSynchronize());
232 }
233 eExecutionPolicy executionPolicy() const override
234 {
236 }
237 impl::IRunQueueStream* createStream(const RunQueueBuildInfo& bi) override
238 {
239 return new HipRunQueueStream(this, bi);
240 }
241 impl::IRunQueueEventImpl* createEventImpl() override
242 {
243 return new HipRunQueueEvent(false);
244 }
245 impl::IRunQueueEventImpl* createEventImplWithTimer() override
246 {
247 return new HipRunQueueEvent(true);
248 }
249 void setMemoryAdvice(ConstMemoryView buffer, eMemoryAdvice advice, DeviceId device_id) override
250 {
251 auto v = buffer.bytes();
252 const void* ptr = v.data();
253 size_t count = v.size();
254 int device = device_id.asInt32();
255 hipMemoryAdvise hip_advise;
256
257 if (advice == eMemoryAdvice::MostlyRead)
258 hip_advise = hipMemAdviseSetReadMostly;
260 hip_advise = hipMemAdviseSetPreferredLocation;
261 else if (advice == eMemoryAdvice::AccessedByDevice)
262 hip_advise = hipMemAdviseSetAccessedBy;
263 else if (advice == eMemoryAdvice::PreferredLocationHost) {
264 hip_advise = hipMemAdviseSetPreferredLocation;
265 device = hipCpuDeviceId;
266 }
267 else if (advice == eMemoryAdvice::AccessedByHost) {
268 hip_advise = hipMemAdviseSetAccessedBy;
269 device = hipCpuDeviceId;
270 }
271 else
272 return;
273 //std::cout << "MEMADVISE p=" << ptr << " size=" << count << " advise = " << hip_advise << " id = " << device << "\n";
274 ARCANE_CHECK_HIP(hipMemAdvise(ptr, count, hip_advise, device));
275 }
276 void unsetMemoryAdvice(ConstMemoryView buffer, eMemoryAdvice advice, DeviceId device_id) override
277 {
278 auto v = buffer.bytes();
279 const void* ptr = v.data();
280 size_t count = v.size();
281 int device = device_id.asInt32();
282 hipMemoryAdvise hip_advise;
283
284 if (advice == eMemoryAdvice::MostlyRead)
285 hip_advise = hipMemAdviseUnsetReadMostly;
287 hip_advise = hipMemAdviseUnsetPreferredLocation;
288 else if (advice == eMemoryAdvice::AccessedByDevice)
289 hip_advise = hipMemAdviseUnsetAccessedBy;
290 else if (advice == eMemoryAdvice::PreferredLocationHost) {
291 hip_advise = hipMemAdviseUnsetPreferredLocation;
292 device = hipCpuDeviceId;
293 }
294 else if (advice == eMemoryAdvice::AccessedByHost) {
295 hip_advise = hipMemAdviseUnsetAccessedBy;
296 device = hipCpuDeviceId;
297 }
298 else
299 return;
300 ARCANE_CHECK_HIP(hipMemAdvise(ptr, count, hip_advise, device));
301 }
302
303 void setCurrentDevice(DeviceId device_id) final
304 {
305 Int32 id = device_id.asInt32();
306 if (!device_id.isAccelerator())
307 ARCANE_FATAL("Device {0} is not an accelerator device", id);
308 ARCANE_CHECK_HIP(hipSetDevice(id));
309 }
310 const IDeviceInfoList* deviceInfoList() override { return &m_device_info_list; }
311
312 void getPointerAttribute(PointerAttribute& attribute, const void* ptr) override
313 {
314 hipPointerAttribute_t pa;
315 hipError_t ret_value = hipPointerGetAttributes(&pa, ptr);
316 auto mem_type = ePointerMemoryType::Unregistered;
317 // Si \a ptr n'a pas été alloué dynamiquement (i.e: il est sur la pile),
318 // hipPointerGetAttribute() retourne une erreur. Dans ce cas on considère
319 // la mémoire comme non enregistrée.
320 if (ret_value==hipSuccess){
321#if HIP_VERSION_MAJOR >= 6
322 auto rocm_memory_type = pa.type;
323#else
324 auto rocm_memory_type = pa.memoryType;
325#endif
326 if (pa.isManaged)
327 mem_type = ePointerMemoryType::Managed;
328 else if (rocm_memory_type == hipMemoryTypeHost)
329 mem_type = ePointerMemoryType::Host;
330 else if (rocm_memory_type == hipMemoryTypeDevice)
331 mem_type = ePointerMemoryType::Device;
332 }
333
334 //std::cout << "HIP Info: hip_memory_type=" << (int)pa.memoryType << " is_managed?=" << pa.isManaged
335 // << " flags=" << pa.allocationFlags
336 // << " my_memory_type=" << (int)mem_type
337 // << "\n";
338 _fillPointerAttribute(attribute, mem_type, pa.device,
339 ptr, pa.devicePointer, pa.hostPointer);
340 }
341
342 DeviceMemoryInfo getDeviceMemoryInfo(DeviceId device_id) override
343 {
344 int d = 0;
345 int wanted_d = device_id.asInt32();
346 ARCANE_CHECK_HIP(hipGetDevice(&d));
347 if (d != wanted_d)
348 ARCANE_CHECK_HIP(hipSetDevice(wanted_d));
349 size_t free_mem = 0;
350 size_t total_mem = 0;
351 ARCANE_CHECK_HIP(hipMemGetInfo(&free_mem, &total_mem));
352 if (d != wanted_d)
353 ARCANE_CHECK_HIP(hipSetDevice(d));
355 dmi.setFreeMemory(free_mem);
356 dmi.setTotalMemory(total_mem);
357 return dmi;
358 }
359
360 void pushProfilerRange(const String& name, [[maybe_unused]] Int32 color) override
361 {
362#ifdef ARCANE_HAS_ROCTX
363 roctxRangePush(name.localstr());
364#endif
365 }
366 void popProfilerRange() override
367 {
368#ifdef ARCANE_HAS_ROCTX
369 roctxRangePop();
370#endif
371 }
372
373 public:
374
375 void fillDevices(bool is_verbose);
376
377 private:
378
379 Int64 m_nb_kernel_launched = 0;
380 bool m_is_verbose = false;
381 impl::DeviceInfoList m_device_info_list;
382};
383
384/*---------------------------------------------------------------------------*/
385/*---------------------------------------------------------------------------*/
386
387void HipRunnerRuntime::
388fillDevices(bool is_verbose)
389{
390 int nb_device = 0;
391 ARCANE_CHECK_HIP(hipGetDeviceCount(&nb_device));
392 std::ostream& omain = std::cout;
393 if (is_verbose)
394 omain << "ArcaneHIP: Initialize Arcane HIP runtime nb_available_device=" << nb_device << "\n";
395 for (int i = 0; i < nb_device; ++i) {
396 OStringStream ostr;
397 std::ostream& o = ostr.stream();
398
399 hipDeviceProp_t dp;
400 ARCANE_CHECK_HIP(hipGetDeviceProperties(&dp, i));
401
402 int has_managed_memory = 0;
403 ARCANE_CHECK_HIP(hipDeviceGetAttribute(&has_managed_memory, hipDeviceAttributeManagedMemory, i));
404
405 // Le format des versions dans HIP est:
406 // HIP_VERSION = (HIP_VERSION_MAJOR * 10000000 + HIP_VERSION_MINOR * 100000 + HIP_VERSION_PATCH)
407
408 int runtime_version = 0;
409 ARCANE_CHECK_HIP(hipRuntimeGetVersion(&runtime_version));
410 //runtime_version /= 10000;
411 int runtime_major = runtime_version / 10000000;
412 int runtime_minor = (runtime_version / 100000) % 100;
413
414 int driver_version = 0;
415 ARCANE_CHECK_HIP(hipDriverGetVersion(&driver_version));
416 //driver_version /= 10000;
417 int driver_major = driver_version / 10000000;
418 int driver_minor = (driver_version / 100000) % 100;
419
420 o << "\nDevice " << i << " name=" << dp.name << "\n";
421 o << " Driver version = " << driver_major << "." << (driver_minor) << "." << (driver_version % 100000) << "\n";
422 o << " Runtime version = " << runtime_major << "." << (runtime_minor) << "." << (runtime_version % 100000) << "\n";
423 o << " computeCapability = " << dp.major << "." << dp.minor << "\n";
424 o << " totalGlobalMem = " << dp.totalGlobalMem << "\n";
425 o << " sharedMemPerBlock = " << dp.sharedMemPerBlock << "\n";
426 o << " regsPerBlock = " << dp.regsPerBlock << "\n";
427 o << " warpSize = " << dp.warpSize << "\n";
428 o << " memPitch = " << dp.memPitch << "\n";
429 o << " maxThreadsPerBlock = " << dp.maxThreadsPerBlock << "\n";
430 o << " totalConstMem = " << dp.totalConstMem << "\n";
431 o << " clockRate = " << dp.clockRate << "\n";
432 //o << " deviceOverlap = " << dp.deviceOverlap<< "\n";
433 o << " multiProcessorCount = " << dp.multiProcessorCount << "\n";
434 o << " kernelExecTimeoutEnabled = " << dp.kernelExecTimeoutEnabled << "\n";
435 o << " integrated = " << dp.integrated << "\n";
436 o << " canMapHostMemory = " << dp.canMapHostMemory << "\n";
437 o << " computeMode = " << dp.computeMode << "\n";
438 o << " maxThreadsDim = " << dp.maxThreadsDim[0] << " " << dp.maxThreadsDim[1]
439 << " " << dp.maxThreadsDim[2] << "\n";
440 o << " maxGridSize = " << dp.maxGridSize[0] << " " << dp.maxGridSize[1]
441 << " " << dp.maxGridSize[2] << "\n";
442 o << " concurrentManagedAccess = " << dp.concurrentManagedAccess << "\n";
443 o << " directManagedMemAccessFromHost = " << dp.directManagedMemAccessFromHost << "\n";
444 o << " gcnArchName = " << dp.gcnArchName << "\n";
445 o << " pageableMemoryAccess = " << dp.pageableMemoryAccess << "\n";
446 o << " pageableMemoryAccessUsesHostPageTables = " << dp.pageableMemoryAccessUsesHostPageTables << "\n";
447 o << " hasManagedMemory = " << has_managed_memory << "\n";
448#if HIP_VERSION_MAJOR >= 6
449 o << " gpuDirectRDMASupported = " << dp.gpuDirectRDMASupported << "\n";
450 o << " hostNativeAtomicSupported = " << dp.hostNativeAtomicSupported << "\n";
451 o << " unifiedFunctionPointers = " << dp.unifiedFunctionPointers << "\n";
452#endif
453 {
454 hipDevice_t device;
455 ARCANE_CHECK_HIP(hipDeviceGet(&device, i));
456 hipUUID device_uuid;
457 ARCANE_CHECK_HIP(hipDeviceGetUuid(&device_uuid, device));
458 o << " deviceUuid=";
459 impl::printUUID(o, device_uuid.bytes);
460 o << "\n";
461 }
462
463 String description(ostr.str());
464 if (is_verbose)
465 omain << description;
466
467 DeviceInfo device_info;
468 device_info.setDescription(description);
469 device_info.setDeviceId(DeviceId(i));
470 device_info.setName(dp.name);
471 device_info.setWarpSize(dp.warpSize);
472 m_device_info_list.addDevice(device_info);
473 }
474}
475
476/*---------------------------------------------------------------------------*/
477/*---------------------------------------------------------------------------*/
478
480: public IMemoryCopier
481{
482 void copy(ConstMemoryView from, [[maybe_unused]] eMemoryRessource from_mem,
483 MutableMemoryView to, [[maybe_unused]] eMemoryRessource to_mem,
484 const RunQueue* queue) override
485 {
486 if (queue) {
487 queue->copyMemory(MemoryCopyArgs(to.bytes(), from.bytes()).addAsync(queue->isAsync()));
488 return;
489 }
490 // 'hipMemcpyDefault' sait automatiquement ce qu'il faut faire en tenant
491 // uniquement compte de la valeur des pointeurs. Il faudrait voir si
492 // utiliser \a from_mem et \a to_mem peut améliorer les performances.
493 ARCANE_CHECK_HIP(hipMemcpy(to.data(), from.data(), from.bytes().size(), hipMemcpyDefault));
494 }
495};
496
497/*---------------------------------------------------------------------------*/
498/*---------------------------------------------------------------------------*/
499
500} // End namespace Arcane::Accelerator::Hip
501
502namespace
503{
505Arcane::Accelerator::Hip::HipMemoryCopier global_hip_memory_copier;
506}
507
508/*---------------------------------------------------------------------------*/
509/*---------------------------------------------------------------------------*/
510
511// Cette fonction est le point d'entrée utilisé lors du chargement
512// dynamique de cette bibliothèque
513extern "C" ARCANE_EXPORT void
514arcaneRegisterAcceleratorRuntimehip(Arcane::Accelerator::RegisterRuntimeInfo& init_info)
515{
516 using namespace Arcane;
517 using namespace Arcane::Accelerator::Hip;
518 Arcane::Accelerator::impl::setUsingHIPRuntime(true);
519 Arcane::Accelerator::impl::setHIPRunQueueRuntime(&global_hip_runtime);
521 MemoryUtils::setAcceleratorHostMemoryAllocator(getHipMemoryAllocator());
523 mrm->setIsAccelerator(true);
524 mrm->setAllocator(eMemoryRessource::UnifiedMemory, getHipUnifiedMemoryAllocator());
525 mrm->setAllocator(eMemoryRessource::HostPinned, getHipHostPinnedMemoryAllocator());
526 mrm->setAllocator(eMemoryRessource::Device, getHipDeviceMemoryAllocator());
527 mrm->setCopier(&global_hip_memory_copier);
528 global_hip_runtime.fillDevices(init_info.isVerbose());
529}
530
531/*---------------------------------------------------------------------------*/
532/*---------------------------------------------------------------------------*/
#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.
Identifiant d'un composant du système.
Definition DeviceId.h:33
bool isHost() const
Indique si l'instance est associée à l'hôte.
Definition DeviceId.h:60
Int32 asInt32() const
Valeur numérique du device.
Definition DeviceId.h:69
bool isAccelerator() const
Indique si l'instance est associée à un accélérateur.
Definition DeviceId.h:66
Information sur un device.
Definition DeviceInfo.h:32
Information mémoire d'un accélérateur.
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.
void notifyBeginLaunchKernel(impl::RunCommandImpl &c) override
Notification avant le lancement de la commande.
bool _barrierNoException() override
Barrière sans exception. Retourne true en cas d'erreur.
void barrier() override
Bloque jusqu'à ce que toutes les actions associées à cette file soient terminées.
void notifyEndLaunchKernel(impl::RunCommandImpl &) override
Notification de fin de lancement de la 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.
impl::NativeStream nativeStream() override
Pointeur sur la structure interne dépendante de l'implémentation.
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
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.
Vue constante sur une zone mémoire contigue contenant des éléments de taille fixe.
constexpr SpanType bytes() const
Vue sous forme d'octets.
constexpr const std::byte * data() const
Pointeur sur la zone mémoire.
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.
Vue modifiable sur une zone mémoire contigue contenant des éléments de taille fixe.
constexpr std::byte * data() const
Pointeur sur la zone mémoire.
constexpr SpanType bytes() const
Vue sous forme d'octets.
Flot de sortie lié à une String.
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:516
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:304
Chaîne de caractères unicode.
const char * localstr() const
Retourne la conversion de l'instance dans l'encodage UTF-8.
Definition String.cc:228
eMemoryAdvice
Conseils pour la gestion mémoire.
Definition Memory.h:36
@ AccessedByHost
Indique que la zone mémoire est accédée par l'hôte.
Definition Memory.h:48
@ PreferredLocationDevice
Privilégié le positionnement de la mémoire sur l'accélérateur.
Definition Memory.h:42
@ MostlyRead
Indique que la zone mémoire est principalement en lecture seule.
Definition Memory.h:40
@ PreferredLocationHost
Privilégié le positionnement de la mémoire sur l'hôte.
Definition Memory.h:44
@ AccessedByDevice
Indique que la zone mémoire est accédée par l'accélérateur.
Definition Memory.h:46
eExecutionPolicy
Politique d'exécution pour un Runner.
@ HIP
Politique d'exécution utilisant l'environnement HIP.
IMemoryRessourceMng * getDataMemoryResourceMng()
Gestionnaire de ressource mémoire pour les données.
IMemoryAllocator * setAcceleratorHostMemoryAllocator(IMemoryAllocator *a)
Positionne l'allocateur spécifique pour les accélérateurs.
void setDefaultDataMemoryResource(eMemoryResource mem_resource)
Positionne la ressource mémoire utilisée pour l'allocateur mémoire des données.
-*- 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.
Espace de nom de Arccore.
Arcane::eMemoryResource eMemoryRessource
Typedef pour la version Arcane historique (avec 2's')