Arcane  v3.16.10.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.
Definition MemoryView.h:38
constexpr SpanType bytes() const
Vue sous forme d'octets.
Definition MemoryView.h:107
constexpr const std::byte * data() const
Pointeur sur la zone mémoire.
Definition MemoryView.h:110
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.
Definition MemoryView.h:156
constexpr std::byte * data() const
Pointeur sur la zone mémoire.
Definition MemoryView.h:218
constexpr SpanType bytes() const
Vue sous forme d'octets.
Definition MemoryView.h:215
Flot de sortie lié à une String.
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:212
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:422
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')