14#include "arcane/accelerator/hip/HipAccelerator.h"
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"
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"
40#ifdef ARCANE_HAS_ROCTX
46namespace Arcane::Accelerator::Hip
52class HipRunQueueStream
61 ARCANE_CHECK_HIP(hipStreamCreate(&m_hip_stream));
63 int priority = bi.priority();
64 ARCANE_CHECK_HIP(hipStreamCreateWithPriority(&m_hip_stream, hipStreamDefault, priority));
67 ~HipRunQueueStream()
override
69 ARCANE_CHECK_HIP_NOTHROW(hipStreamDestroy(m_hip_stream));
76#ifdef ARCANE_HAS_ROCTX
77 auto kname = c.kernelName();
79 roctxRangePush(c.traceInfo().name());
81 roctxRangePush(kname.localstr());
83 return m_runtime->notifyBeginLaunchKernel();
87#ifdef ARCANE_HAS_ROCTX
90 return m_runtime->notifyEndLaunchKernel();
94 ARCANE_CHECK_HIP(hipStreamSynchronize(m_hip_stream));
98 return hipStreamSynchronize(m_hip_stream) != hipSuccess;
102 auto r = hipMemcpyAsync(args.destination().
data(), args.source().
data(),
103 args.source().
bytes().
size(), hipMemcpyDefault, m_hip_stream);
110 auto src = args.source().
bytes();
114 int device = hipCpuDeviceId;
117 auto r = hipMemPrefetchAsync(src.data(), src.size(), device, m_hip_stream);
129 hipStream_t trueStream()
const
137 hipStream_t m_hip_stream;
143class HipRunQueueEvent
148 explicit HipRunQueueEvent(
bool has_timer)
151 ARCANE_CHECK_HIP(hipEventCreate(&m_hip_event));
153 ARCANE_CHECK_HIP(hipEventCreateWithFlags(&m_hip_event, hipEventDisableTiming));
155 ~HipRunQueueEvent()
override
157 ARCANE_CHECK_HIP_NOTHROW(hipEventDestroy(m_hip_event));
166 ARCANE_CHECK_HIP(hipEventRecord(m_hip_event, rq->trueStream()));
171 ARCANE_CHECK_HIP(hipEventSynchronize(m_hip_event));
177 ARCANE_CHECK_HIP(hipStreamWaitEvent(rq->trueStream(), m_hip_event, 0));
180 Int64 elapsedTime(IRunQueueEventImpl* from_event)
final
182 auto* true_from_event =
static_cast<HipRunQueueEvent*
>(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);
191 bool hasPendingWork()
final
193 hipError_t v = hipEventQuery(m_hip_event);
194 if (v == hipErrorNotReady)
202 hipEvent_t m_hip_event;
217 void notifyBeginLaunchKernel()
override
219 ++m_nb_kernel_launched;
221 std::cout <<
"BEGIN HIP KERNEL!\n";
223 void notifyEndLaunchKernel()
override
225 ARCANE_CHECK_HIP(hipGetLastError());
227 std::cout <<
"END HIP KERNEL!\n";
229 void barrier()
override
231 ARCANE_CHECK_HIP(hipDeviceSynchronize());
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;
258 hip_advise = hipMemAdviseSetReadMostly;
260 hip_advise = hipMemAdviseSetPreferredLocation;
262 hip_advise = hipMemAdviseSetAccessedBy;
264 hip_advise = hipMemAdviseSetPreferredLocation;
265 device = hipCpuDeviceId;
268 hip_advise = hipMemAdviseSetAccessedBy;
269 device = hipCpuDeviceId;
274 ARCANE_CHECK_HIP(hipMemAdvise(ptr, count, hip_advise, device));
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;
285 hip_advise = hipMemAdviseUnsetReadMostly;
287 hip_advise = hipMemAdviseUnsetPreferredLocation;
289 hip_advise = hipMemAdviseUnsetAccessedBy;
291 hip_advise = hipMemAdviseUnsetPreferredLocation;
292 device = hipCpuDeviceId;
295 hip_advise = hipMemAdviseUnsetAccessedBy;
296 device = hipCpuDeviceId;
300 ARCANE_CHECK_HIP(hipMemAdvise(ptr, count, hip_advise, device));
303 void setCurrentDevice(
DeviceId device_id)
final
305 Int32
id = device_id.
asInt32();
307 ARCANE_FATAL(
"Device {0} is not an accelerator device",
id);
308 ARCANE_CHECK_HIP(hipSetDevice(
id));
310 const IDeviceInfoList* deviceInfoList()
override {
return &m_device_info_list; }
312 void getPointerAttribute(
PointerAttribute& attribute,
const void* ptr)
override
314 hipPointerAttribute_t pa;
315 hipError_t ret_value = hipPointerGetAttributes(&pa, ptr);
316 auto mem_type = ePointerMemoryType::Unregistered;
320 if (ret_value==hipSuccess){
321#if HIP_VERSION_MAJOR >= 6
322 auto rocm_memory_type = pa.type;
324 auto rocm_memory_type = pa.memoryType;
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;
338 _fillPointerAttribute(attribute, mem_type, pa.device,
339 ptr, pa.devicePointer, pa.hostPointer);
345 int wanted_d = device_id.
asInt32();
346 ARCANE_CHECK_HIP(hipGetDevice(&d));
348 ARCANE_CHECK_HIP(hipSetDevice(wanted_d));
350 size_t total_mem = 0;
351 ARCANE_CHECK_HIP(hipMemGetInfo(&free_mem, &total_mem));
353 ARCANE_CHECK_HIP(hipSetDevice(d));
355 dmi.setFreeMemory(free_mem);
356 dmi.setTotalMemory(total_mem);
360 void pushProfilerRange(
const String& name, [[maybe_unused]] Int32 color)
override
362#ifdef ARCANE_HAS_ROCTX
366 void popProfilerRange()
override
368#ifdef ARCANE_HAS_ROCTX
375 void fillDevices(
bool is_verbose);
379 Int64 m_nb_kernel_launched = 0;
380 bool m_is_verbose =
false;
387void HipRunnerRuntime::
388fillDevices(
bool is_verbose)
391 ARCANE_CHECK_HIP(hipGetDeviceCount(&nb_device));
392 std::ostream& omain = std::cout;
394 omain <<
"ArcaneHIP: Initialize Arcane HIP runtime nb_available_device=" << nb_device <<
"\n";
395 for (
int i = 0; i < nb_device; ++i) {
397 std::ostream& o = ostr.stream();
400 ARCANE_CHECK_HIP(hipGetDeviceProperties(&dp, i));
402 int has_managed_memory = 0;
403 ARCANE_CHECK_HIP(hipDeviceGetAttribute(&has_managed_memory, hipDeviceAttributeManagedMemory, i));
408 int runtime_version = 0;
409 ARCANE_CHECK_HIP(hipRuntimeGetVersion(&runtime_version));
411 int runtime_major = runtime_version / 10000000;
412 int runtime_minor = (runtime_version / 100000) % 100;
414 int driver_version = 0;
415 ARCANE_CHECK_HIP(hipDriverGetVersion(&driver_version));
417 int driver_major = driver_version / 10000000;
418 int driver_minor = (driver_version / 100000) % 100;
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";
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";
455 ARCANE_CHECK_HIP(hipDeviceGet(&device, i));
457 ARCANE_CHECK_HIP(hipDeviceGetUuid(&device_uuid, device));
459 impl::printUUID(o, device_uuid.bytes);
463 String description(ostr.str());
465 omain << description;
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);
493 ARCANE_CHECK_HIP(hipMemcpy(to.
data(), from.
data(), from.
bytes().
size(), hipMemcpyDefault));
513extern "C" ARCANE_EXPORT
void
517 using namespace Arcane::Accelerator::Hip;
518 Arcane::Accelerator::impl::setUsingHIPRuntime(
true);
519 Arcane::Accelerator::impl::setHIPRunQueueRuntime(&global_hip_runtime);
527 mrm->
setCopier(&global_hip_memory_copier);
528 global_hip_runtime.fillDevices(init_info.isVerbose());
#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.
bool isHost() const
Indique si l'instance est associée à l'hôte.
Int32 asInt32() const
Valeur numérique du device.
bool isAccelerator() const
Indique si l'instance est associée à un accélérateur.
Information sur un device.
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.
Arguments pour le préfetching mémoire.
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.
void copyMemory(const MemoryCopyArgs &args) const
Copie des informations entre deux zones mémoires.
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__ SizeType size() const noexcept
Retourne la taille du tableau.
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Chaîne de caractères unicode.
const char * localstr() const
Retourne la conversion de l'instance dans l'encodage UTF-8.
eMemoryAdvice
Conseils pour la gestion mémoire.
@ AccessedByHost
Indique que la zone mémoire est accédée par l'hôte.
@ PreferredLocationDevice
Privilégié le positionnement de la mémoire sur l'accélérateur.
@ MostlyRead
Indique que la zone mémoire est principalement en lecture seule.
@ PreferredLocationHost
Privilégié le positionnement de la mémoire sur l'hôte.
@ AccessedByDevice
Indique que la zone mémoire est accédée par l'accélérateur.
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')