1//===- InstrProfilingPlatformROCm.cpp - Profile data ROCm platform -------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8
9extern "C" {
10#include "InstrProfiling.h"
11#include "InstrProfilingPort.h"
12#include "InstrProfilingUtil.h"
13}
14
15#include "interception/interception.h"
16// C library headers (not <cstdio> etc.): clang_rt.profile is built with
17// -nostdinc++ and avoids the C++ standard library (see profile/CMakeLists.txt).
18#include <stddef.h>
19#include <stdio.h>
20#include <stdlib.h>
21#include <string.h>
22#ifdef _WIN32
23#include <wchar.h>
24#endif
25
26#ifdef _WIN32
27#define WIN32_LEAN_AND_MEAN
28#include <windows.h>
29// windows.h needs to be included before tlhelp32.h.
30#include <tlhelp32.h>
31#else
32#include <dlfcn.h>
33#include <pthread.h>
34#endif
35
36#include "InstrProfilingPlatformROCmInternal.h"
37
38// shortcut to shared helper names
39using namespace __prof_rocm;
40
41/* Serialize one-time HIP loader resolution and DynamicModules mutations.
42 * Inline to avoid a sanitizer_common dependency. */
43#ifdef _WIN32
44static INIT_ONCE HipLoadedOnce = INIT_ONCE_STATIC_INIT;
45static CRITICAL_SECTION DynamicModulesLock;
46static INIT_ONCE DynamicModulesLockInit = INIT_ONCE_STATIC_INIT;
47static BOOL CALLBACK initDynamicModulesLockCb(PINIT_ONCE, PVOID, PVOID *) {
48 InitializeCriticalSection(&DynamicModulesLock);
49 return TRUE;
50}
51static void lockDynamicModules(void) {
52 InitOnceExecuteOnce(&DynamicModulesLockInit, initDynamicModulesLockCb, NULL,
53 NULL);
54 EnterCriticalSection(&DynamicModulesLock);
55}
56static void unlockDynamicModules(void) {
57 LeaveCriticalSection(&DynamicModulesLock);
58}
59#else
60static pthread_once_t HipLoadedOnce = PTHREAD_ONCE_INIT;
61static pthread_mutex_t DynamicModulesLock = PTHREAD_MUTEX_INITIALIZER;
62static void lockDynamicModules(void) {
63 pthread_mutex_lock(mutex: &DynamicModulesLock);
64}
65static void unlockDynamicModules(void) {
66 pthread_mutex_unlock(mutex: &DynamicModulesLock);
67}
68#endif
69
70int __prof_rocm::isVerboseMode() {
71 static int IsVerbose = -1;
72 if (IsVerbose == -1)
73 IsVerbose = getenv(name: "LLVM_PROFILE_VERBOSE") != nullptr;
74 return IsVerbose;
75}
76
77/* -------------------------------------------------------------------------- */
78/* Dynamic loading of HIP runtime symbols */
79/* -------------------------------------------------------------------------- */
80
81typedef int (*hipGetSymbolAddressTy)(void **, const void *);
82typedef int (*hipGetSymbolSizeTy)(size_t *, const void *);
83typedef int (*hipMemcpyTy)(void *, const void *, size_t, int);
84typedef int (*hipModuleGetGlobalTy)(void **, size_t *, void *, const char *);
85typedef int (*hipGetDeviceCountTy)(int *);
86typedef int (*hipGetDeviceTy)(int *);
87typedef int (*hipSetDeviceTy)(int);
88#if defined(__linux__) && !defined(_WIN32)
89typedef void *HipStream;
90typedef int (*hipStreamGetDeviceTy)(HipStream, int *);
91#endif
92
93/* Minimal hipDeviceProp_t (HIP 6.x R0600): only gcnArchName at offset 1160
94 * is read. Padded to 4096 to tolerate ABI growth. */
95typedef struct {
96 char padding[1160];
97 char gcnArchName[256];
98 char tail_padding[2680];
99} HipDevicePropMinimal;
100typedef int (*hipGetDevicePropertiesTy)(HipDevicePropMinimal *, int);
101
102static hipGetSymbolAddressTy pHipGetSymbolAddress = nullptr;
103static hipGetSymbolSizeTy pHipGetSymbolSize = nullptr;
104static hipMemcpyTy pHipMemcpy = nullptr;
105static hipModuleGetGlobalTy pHipModuleGetGlobal = nullptr;
106static hipGetDeviceCountTy pHipGetDeviceCount = nullptr;
107static hipGetDeviceTy pHipGetDevice = nullptr;
108static hipSetDeviceTy pHipSetDevice = nullptr;
109#if defined(__linux__) && !defined(_WIN32)
110static hipStreamGetDeviceTy pHipStreamGetDevice = nullptr;
111#endif
112static hipGetDevicePropertiesTy pHipGetDeviceProperties = nullptr;
113
114static int NumDevices = 0;
115/* 256 matches hipDeviceProp_t::gcnArchName, the source field width. */
116static char (*DeviceArchNames)[256] = nullptr;
117#if defined(__linux__) && !defined(_WIN32)
118static unsigned char *UsedDevices = nullptr;
119static int AnyDeviceUsed = 0;
120#endif
121
122#ifdef _WIN32
123static wchar_t toLowerAsciiW(wchar_t C) {
124 return C >= L'A' && C <= L'Z' ? C - L'A' + L'a' : C;
125}
126
127static int wcsEqualNoCase(const wchar_t *A, const wchar_t *B) {
128 while (*A && *B) {
129 if (toLowerAsciiW(*A) != toLowerAsciiW(*B))
130 return 0;
131 ++A;
132 ++B;
133 }
134 return *A == *B;
135}
136
137static int wcsStartsWithNoCase(const wchar_t *S, const wchar_t *Prefix) {
138 while (*Prefix) {
139 if (toLowerAsciiW(*S) != toLowerAsciiW(*Prefix))
140 return 0;
141 ++S;
142 ++Prefix;
143 }
144 return 1;
145}
146
147static int wcsEndsWithNoCase(const wchar_t *S, const wchar_t *Suffix) {
148 size_t SLen = wcslen(S);
149 size_t SuffixLen = wcslen(Suffix);
150 return SLen >= SuffixLen && wcsEqualNoCase(S + SLen - SuffixLen, Suffix);
151}
152
153static int isHipRuntimeModuleName(const wchar_t *Name) {
154 return wcsEqualNoCase(Name, L"amdhip64.dll") ||
155 (wcsStartsWithNoCase(Name, L"amdhip64_") &&
156 wcsEndsWithNoCase(Name, L".dll"));
157}
158
159static void *findLoadedHipRuntime(void) {
160 HMODULE Handle = GetModuleHandleW(L"amdhip64.dll");
161 if (Handle)
162 return (void *)Handle;
163
164 HANDLE Snapshot = CreateToolhelp32Snapshot(
165 TH32CS_SNAPMODULE | TH32CS_SNAPMODULE32, GetCurrentProcessId());
166 if (Snapshot == INVALID_HANDLE_VALUE)
167 return nullptr;
168
169 MODULEENTRY32W Entry;
170 Entry.dwSize = sizeof(Entry);
171 if (Module32FirstW(Snapshot, &Entry)) {
172 do {
173 if (isHipRuntimeModuleName(Entry.szModule)) {
174 Handle = Entry.hModule;
175 break;
176 }
177 } while (Module32NextW(Snapshot, &Entry));
178 }
179
180 CloseHandle(Snapshot);
181 return (void *)Handle;
182}
183#endif
184
185/* -------------------------------------------------------------------------- */
186/* Device-to-host copies */
187/* Keep HIP-only to avoid an HSA dependency. */
188/* -------------------------------------------------------------------------- */
189
190static void doEnsureHipLoaded(void) {
191 if (!__interception::DynamicLoaderAvailable()) {
192 if (isVerboseMode())
193 PROF_NOTE("%s", "Dynamic library loading not available - "
194 "HIP profiling disabled\n");
195 return;
196 }
197
198#ifdef _WIN32
199 /* Use the app's loaded HIP runtime to avoid binding another ROCm version. */
200 void *Handle = findLoadedHipRuntime();
201#else
202 const char *HipLibName = "libamdhip64.so";
203 void *Handle = __interception::OpenLibrary(name: HipLibName);
204#endif
205 if (!Handle)
206 return;
207
208 pHipGetSymbolAddress = (hipGetSymbolAddressTy)__interception::LookupSymbol(
209 handle: Handle, symbol: "hipGetSymbolAddress");
210 pHipGetSymbolSize = (hipGetSymbolSizeTy)__interception::LookupSymbol(
211 handle: Handle, symbol: "hipGetSymbolSize");
212 pHipMemcpy = (hipMemcpyTy)__interception::LookupSymbol(handle: Handle, symbol: "hipMemcpy");
213 pHipModuleGetGlobal = (hipModuleGetGlobalTy)__interception::LookupSymbol(
214 handle: Handle, symbol: "hipModuleGetGlobal");
215 pHipGetDeviceCount = (hipGetDeviceCountTy)__interception::LookupSymbol(
216 handle: Handle, symbol: "hipGetDeviceCount");
217 pHipGetDevice =
218 (hipGetDeviceTy)__interception::LookupSymbol(handle: Handle, symbol: "hipGetDevice");
219 pHipSetDevice =
220 (hipSetDeviceTy)__interception::LookupSymbol(handle: Handle, symbol: "hipSetDevice");
221#if defined(__linux__) && !defined(_WIN32)
222 pHipStreamGetDevice = (hipStreamGetDeviceTy)__interception::LookupSymbol(
223 handle: Handle, symbol: "hipStreamGetDevice");
224#endif
225 pHipGetDeviceProperties =
226 (hipGetDevicePropertiesTy)__interception::LookupSymbol(
227 handle: Handle, symbol: "hipGetDevicePropertiesR0600");
228 if (!pHipGetDeviceProperties)
229 pHipGetDeviceProperties =
230 (hipGetDevicePropertiesTy)__interception::LookupSymbol(
231 handle: Handle, symbol: "hipGetDeviceProperties");
232
233 if (pHipGetDeviceCount && pHipGetDeviceProperties) {
234 int Count = 0;
235 if (pHipGetDeviceCount(&Count) == 0 && Count > 0) {
236 DeviceArchNames = (char (*)[256])calloc(nmemb: Count, size: sizeof(*DeviceArchNames));
237 if (!DeviceArchNames) {
238 PROF_ERR("%s\n", "failed to allocate device arch name table");
239 return;
240 }
241#if defined(__linux__) && !defined(_WIN32)
242 UsedDevices = (unsigned char *)calloc(nmemb: Count, size: sizeof(*UsedDevices));
243 if (!UsedDevices && isVerboseMode())
244 PROF_NOTE("%s\n", "Device-use tracking disabled");
245#endif
246 HipDevicePropMinimal Prop;
247 for (int i = 0; i < Count; ++i) {
248 __builtin_memset(&Prop, 0, sizeof(Prop));
249 if (pHipGetDeviceProperties(&Prop, i) == 0) {
250 strncpy(dest: DeviceArchNames[i], src: Prop.gcnArchName,
251 n: sizeof(DeviceArchNames[i]) - 1);
252 DeviceArchNames[i][sizeof(DeviceArchNames[i]) - 1] = '\0';
253 if (isVerboseMode())
254 PROF_NOTE("Device %d arch: %s\n", i, DeviceArchNames[i]);
255 }
256 }
257 NumDevices = Count;
258 }
259 }
260}
261
262#ifdef _WIN32
263static BOOL CALLBACK ensureHipLoadedCb(PINIT_ONCE, PVOID, PVOID *) {
264 doEnsureHipLoaded();
265 return TRUE;
266}
267#endif
268
269void __prof_rocm::ensureHipLoaded(void) {
270#ifdef _WIN32
271 InitOnceExecuteOnce(&HipLoadedOnce, ensureHipLoadedCb, NULL, NULL);
272#else
273 pthread_once(once_control: &HipLoadedOnce, init_routine: doEnsureHipLoaded);
274#endif
275}
276
277// Accessor for the HSA drain: true once the loaded HIP runtime exposes
278// hipMemcpy. Kept here so pHipMemcpy stays file-private to this TU.
279int __prof_rocm::hipMemcpyAvailable() { return pHipMemcpy != nullptr; }
280
281/* -------------------------------------------------------------------------- */
282/* Public wrappers that forward to the loaded HIP symbols */
283/* -------------------------------------------------------------------------- */
284
285static int hipGetSymbolAddress(void **devPtr, const void *symbol) {
286 ensureHipLoaded();
287 return pHipGetSymbolAddress ? pHipGetSymbolAddress(devPtr, symbol) : -1;
288}
289
290static int hipGetSymbolSize(size_t *size, const void *symbol) {
291 ensureHipLoaded();
292 return pHipGetSymbolSize ? pHipGetSymbolSize(size, symbol) : -1;
293}
294
295static int hipMemcpy(void *dest, const void *src, size_t len,
296 int kind /*2=DToH*/) {
297 ensureHipLoaded();
298 return pHipMemcpy ? pHipMemcpy(dest, src, len, kind) : -1;
299}
300
301/* Device section symbols must be registered with CLR first; otherwise
302 * hipMemcpy may take a CPU path and crash. */
303int __prof_rocm::memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) {
304 return hipMemcpy(dest: Dst, src: Src, len: Size, kind: 2 /* DToH */);
305}
306
307[[maybe_unused]]
308static int hipModuleGetGlobal(void **DevPtr, size_t *Bytes, void *Module,
309 const char *Name) {
310 ensureHipLoaded();
311 return pHipModuleGetGlobal ? pHipModuleGetGlobal(DevPtr, Bytes, Module, Name)
312 : -1;
313}
314
315static int hipGetDevice(int *DeviceId) {
316 ensureHipLoaded();
317 return pHipGetDevice ? pHipGetDevice(DeviceId) : -1;
318}
319
320static int hipSetDevice(int DeviceId) {
321 ensureHipLoaded();
322 return pHipSetDevice ? pHipSetDevice(DeviceId) : -1;
323}
324
325#if defined(__linux__) && !defined(_WIN32)
326static int hipStreamGetDevice(HipStream Stream, int *DeviceId) {
327 ensureHipLoaded();
328 return pHipStreamGetDevice ? pHipStreamGetDevice(Stream, DeviceId) : -1;
329}
330
331static void markDeviceUsed(int DeviceId) {
332 if (DeviceId < 0 || DeviceId >= NumDevices || !UsedDevices)
333 return;
334 __atomic_store_n(&UsedDevices[DeviceId], 1, __ATOMIC_RELAXED);
335 __atomic_store_n(&AnyDeviceUsed, 1, __ATOMIC_RELEASE);
336}
337
338static void markCurrentDeviceUsed(void) {
339 int DeviceId = -1;
340 if (hipGetDevice(DeviceId: &DeviceId) == 0)
341 markDeviceUsed(DeviceId);
342}
343
344static void markLaunchStreamDeviceUsed(HipStream Stream) {
345 int DeviceId = -1;
346 if (Stream && hipStreamGetDevice(Stream, DeviceId: &DeviceId) == 0) {
347 markDeviceUsed(DeviceId);
348 return;
349 }
350 markCurrentDeviceUsed();
351}
352
353static int shouldCollectDevice(int DeviceId) {
354 if (UsedDevices && __atomic_load_n(&AnyDeviceUsed, __ATOMIC_ACQUIRE) &&
355 !__atomic_load_n(&UsedDevices[DeviceId], __ATOMIC_RELAXED))
356 return 0;
357 return 1;
358}
359#else
360static int shouldCollectDevice(int) { return 1; }
361#endif
362
363static const char *getDeviceArchName(int DeviceId) {
364 if (DeviceId < 0 || DeviceId >= NumDevices || !DeviceArchNames[DeviceId][0])
365 return "amdgpu";
366 return DeviceArchNames[DeviceId];
367}
368
369/* -------------------------------------------------------------------------- */
370/* Dynamic module tracking */
371/* -------------------------------------------------------------------------- */
372
373/* Per-TU profile entry inside a dynamic module.
374 * A single dynamic module may contain multiple TUs (e.g. -fgpu-rdc). */
375typedef struct {
376 void *DeviceVar; /* device address of __llvm_profile_sections_<CUID> */
377 int Processed; /* 0 = not yet collected, 1 = data already copied */
378} OffloadDynamicTUInfo;
379
380/* One entry per hipModuleLoad call. */
381typedef struct {
382 void *ModulePtr; /* hipModule_t handle */
383 OffloadDynamicTUInfo *TUs; /* array of per-TU entries */
384 int NumTUs;
385 int CapTUs;
386} OffloadDynamicModuleInfo;
387
388static OffloadDynamicModuleInfo *DynamicModules = nullptr;
389static int NumDynamicModules = 0;
390static int CapDynamicModules = 0;
391
392/* -------------------------------------------------------------------------- */
393/* ELF symbol enumeration (manual parse: compiler-rt cannot link LLVM Support)
394 */
395/* -------------------------------------------------------------------------- */
396
397#if __has_include(<elf.h>)
398#include <elf.h>
399
400/* Callback invoked for every matching symbol name found in the ELF image.
401 * Return 0 to continue iteration, non-zero to stop. */
402typedef int (*SymbolCallback)(const char *Name, void *UserData);
403
404/* If Image is a clang offload bundle, return a pointer to the first embedded
405 * ELF. Returns Image if not a bundle, nullptr if a bundle holds no ELF. */
406static const void *unwrapOffloadBundle(const void *Image) {
407 static const char BundleMagic[] = "__CLANG_OFFLOAD_BUNDLE__";
408 if (memcmp(s1: Image, s2: BundleMagic, n: sizeof(BundleMagic) - 1) != 0)
409 return Image; /* Not a bundle, return as-is. */
410
411 const char *Buf = (const char *)Image;
412 uint64_t NumEntries;
413 __builtin_memcpy(&NumEntries, Buf + sizeof(BundleMagic) - 1,
414 sizeof(uint64_t));
415
416 /* Walk the entry table (starts at offset 32). */
417 const char *Cursor = Buf + 32;
418 for (uint64_t I = 0; I < NumEntries; ++I) {
419 uint64_t EntryOffset, EntrySize, IDSize;
420 __builtin_memcpy(&EntryOffset, Cursor, sizeof(EntryOffset));
421 Cursor += sizeof(EntryOffset);
422 __builtin_memcpy(&EntrySize, Cursor, sizeof(EntrySize));
423 Cursor += sizeof(EntrySize);
424 __builtin_memcpy(&IDSize, Cursor, sizeof(IDSize));
425 Cursor += sizeof(IDSize);
426 Cursor += IDSize; /* skip entry ID */
427
428 if (EntrySize >= sizeof(Elf64_Ehdr)) {
429 const Elf64_Ehdr *E = (const Elf64_Ehdr *)(Buf + EntryOffset);
430 if (E->e_ident[EI_MAG0] == ELFMAG0 && E->e_ident[EI_MAG1] == ELFMAG1 &&
431 E->e_ident[EI_MAG2] == ELFMAG2 && E->e_ident[EI_MAG3] == ELFMAG3) {
432 return (const void *)(Buf + EntryOffset);
433 }
434 }
435 }
436
437 PROF_WARN("%s", "offload bundle contains no valid ELF entries\n");
438 return nullptr;
439}
440
441/* Invoke CB for every global symbol in Image (an AMDGPU ELF or offload bundle)
442 * whose name starts with PREFIX. Image may be null. */
443static void enumerateElfSymbols(const void *Image, const char *Prefix,
444 SymbolCallback CB, void *UserData) {
445 if (!Image)
446 return;
447
448 Image = unwrapOffloadBundle(Image);
449 if (!Image)
450 return;
451
452 const Elf64_Ehdr *Ehdr = (const Elf64_Ehdr *)Image;
453 if (Ehdr->e_ident[EI_MAG0] != ELFMAG0 || Ehdr->e_ident[EI_MAG1] != ELFMAG1 ||
454 Ehdr->e_ident[EI_MAG2] != ELFMAG2 || Ehdr->e_ident[EI_MAG3] != ELFMAG3) {
455 if (isVerboseMode())
456 PROF_NOTE("%s", "Image is not a valid ELF, skipping enumeration\n");
457 return;
458 }
459
460 size_t PrefixLen = strlen(s: Prefix);
461 const char *Base = (const char *)Image;
462 const Elf64_Shdr *Shdrs = (const Elf64_Shdr *)(Base + Ehdr->e_shoff);
463
464 for (int i = 0; i < Ehdr->e_shnum; ++i) {
465 if (Shdrs[i].sh_type != SHT_SYMTAB)
466 continue;
467
468 const Elf64_Sym *Syms = (const Elf64_Sym *)(Base + Shdrs[i].sh_offset);
469 int NumSyms = Shdrs[i].sh_size / sizeof(Elf64_Sym);
470 /* String table is the section referenced by sh_link. */
471 const char *StrTab = Base + Shdrs[Shdrs[i].sh_link].sh_offset;
472
473 for (int j = 0; j < NumSyms; ++j) {
474 if (Syms[j].st_name == 0)
475 continue;
476 const char *Name = StrTab + Syms[j].st_name;
477 if (strncmp(s1: Name, s2: Prefix, n: PrefixLen) == 0) {
478 if (CB(Name, UserData))
479 return;
480 }
481 }
482 }
483}
484
485/* State passed through the enumeration callback. */
486typedef struct {
487 void *Module; /* hipModule_t */
488 OffloadDynamicModuleInfo *ModInfo;
489} EnumState;
490
491/* Register one __llvm_profile_sections_<CUID> symbol on the module entry.
492 * hipModuleGetGlobal also registers the device address with CLR so hipMemcpy
493 * can copy from it later. */
494static int registerPrfSymbol(const char *Name, void *UserData) {
495 EnumState *S = (EnumState *)UserData;
496 OffloadDynamicModuleInfo *MI = S->ModInfo;
497
498 /* The symbol is the per-TU sections struct itself, not a pointer
499 * indirection, so this address is the hipMemcpy source. */
500 void *DeviceVar = nullptr;
501 size_t Bytes = 0;
502 if (hipModuleGetGlobal(DevPtr: &DeviceVar, Bytes: &Bytes, Module: S->Module, Name) != 0) {
503 PROF_WARN("failed to get symbol %s for module %p\n", Name, S->Module);
504 return 0; /* continue */
505 }
506
507 if (growArray(Arr: (void **)&MI->TUs, Cap: &MI->CapTUs, MinCount: MI->NumTUs + 1, InitCap: 4,
508 ElemSize: sizeof(*MI->TUs))) {
509 PROF_ERR("%s\n", "failed to grow TU array");
510 return 0;
511 }
512 OffloadDynamicTUInfo *TU = &MI->TUs[MI->NumTUs++];
513 TU->DeviceVar = DeviceVar;
514 TU->Processed = 0;
515
516 (void)Name;
517 return 0; /* continue enumeration */
518}
519
520#endif /* __has_include(<elf.h>) */
521
522/* -------------------------------------------------------------------------- */
523/* Registration / un-registration helpers */
524/* -------------------------------------------------------------------------- */
525
526extern "C" void
527__llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
528 const void *Image) {
529 if (ModuleLoadRc)
530 return;
531
532 lockDynamicModules();
533
534 if (isVerboseMode())
535 PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n",
536 NumDynamicModules, ModuleLoadRc, *Ptr, Image);
537
538 if (growArray(Arr: (void **)&DynamicModules, Cap: &CapDynamicModules,
539 MinCount: NumDynamicModules + 1, InitCap: 64, ElemSize: sizeof(*DynamicModules))) {
540 unlockDynamicModules();
541 return;
542 }
543
544 OffloadDynamicModuleInfo *MI = &DynamicModules[NumDynamicModules++];
545 MI->ModulePtr = *Ptr;
546 MI->TUs = nullptr;
547 MI->NumTUs = 0;
548 MI->CapTUs = 0;
549
550 /* Dynamic-module profiling needs ELF parsing for symbol enumeration. */
551#if __has_include(<elf.h>)
552 EnumState State = {.Module: *Ptr, .ModInfo: MI};
553 enumerateElfSymbols(Image, Prefix: "__llvm_profile_sections_", CB: registerPrfSymbol,
554 UserData: &State);
555#else
556 (void)Image;
557 if (isVerboseMode())
558 PROF_NOTE("%s",
559 "Dynamic module profiling not supported on this platform\n");
560#endif
561
562 if (MI->NumTUs == 0) {
563 PROF_WARN("no __llvm_profile_sections_* symbols found in module %p\n",
564 *Ptr);
565 } else if (isVerboseMode()) {
566 PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs);
567 }
568
569 unlockDynamicModules();
570}
571
572extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
573 lockDynamicModules();
574 for (int i = 0; i < NumDynamicModules; ++i) {
575 OffloadDynamicModuleInfo *MI = &DynamicModules[i];
576
577 /* HIP recycles hipModule_t addresses; drained slots are cleared so a
578 * recycled handle finds the new slot, not the dead one. */
579 if (MI->ModulePtr != Ptr)
580 continue;
581
582 if (isVerboseMode())
583 PROF_NOTE("Unregistering module %p (%d TUs)\n", MI->ModulePtr,
584 MI->NumTUs);
585
586 static int NextTUIndex = 0;
587 for (int t = 0; t < MI->NumTUs; ++t) {
588 OffloadDynamicTUInfo *TU = &MI->TUs[t];
589 if (TU->Processed) {
590 if (isVerboseMode())
591 PROF_NOTE("Module %p TU %d already processed, skipping\n", Ptr, t);
592 continue;
593 }
594 int TUIndex = __atomic_fetch_add(&NextTUIndex, 1, __ATOMIC_RELAXED);
595 if (TU->DeviceVar) {
596 int CurDev = 0;
597 hipGetDevice(DeviceId: &CurDev);
598 const char *ArchName = getDeviceArchName(DeviceId: CurDev);
599 /* Encode TUIndex in Target so each drain writes a distinct profraw;
600 * otherwise back-to-back drains overwrite the same file. */
601 char TargetWithTU[64];
602 snprintf(s: TargetWithTU, maxlen: sizeof(TargetWithTU), format: "%s.%d", ArchName,
603 TUIndex);
604 if (processDeviceOffloadPrf(DeviceOffloadPrf: TU->DeviceVar, Target: TargetWithTU, Sections: nullptr) == 0)
605 TU->Processed = 1;
606 else
607 PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr,
608 t);
609 }
610 }
611 MI->ModulePtr = nullptr;
612 unlockDynamicModules();
613 return;
614 }
615
616 if (isVerboseMode())
617 PROF_WARN("unregister called for unknown module %p\n", Ptr);
618 unlockDynamicModules();
619}
620
621static void **OffloadShadowVariables = nullptr;
622static int NumShadowVariables = 0;
623static int CapShadowVariables = 0;
624
625struct OffloadSectionShadow {
626 void *Data;
627 void *Counters;
628 void *UniformCounters;
629 void *Names;
630};
631
632struct OffloadSectionShadowGroup {
633 OffloadSectionShadow *Shadows;
634 int NumShadows;
635 int CapShadows;
636 int NumSections;
637};
638
639static OffloadSectionShadowGroup *OffloadSectionShadowGroups = nullptr;
640static int CapSectionShadowGroups = 0;
641
642static int ensureSectionShadowGroupCapacity(void) {
643 return growArray(Arr: (void **)&OffloadSectionShadowGroups,
644 Cap: &CapSectionShadowGroups, MinCount: CapShadowVariables,
645 InitCap: CapShadowVariables, ElemSize: sizeof(*OffloadSectionShadowGroups));
646}
647
648static int ensureSectionShadowCapacity(OffloadSectionShadowGroup *Group,
649 int MinCapacity) {
650 return growArray(Arr: (void **)&Group->Shadows, Cap: &Group->CapShadows, MinCount: MinCapacity, InitCap: 4,
651 ElemSize: sizeof(*Group->Shadows));
652}
653
654extern "C" void __llvm_profile_offload_register_shadow_variable(void *ptr) {
655 if (growArray(Arr: (void **)&OffloadShadowVariables, Cap: &CapShadowVariables,
656 MinCount: NumShadowVariables + 1, InitCap: 64, ElemSize: sizeof(*OffloadShadowVariables)))
657 return;
658 if (ensureSectionShadowGroupCapacity())
659 return;
660 int Index = NumShadowVariables++;
661 OffloadShadowVariables[Index] = ptr;
662 __builtin_memset(&OffloadSectionShadowGroups[Index], 0,
663 sizeof(OffloadSectionShadowGroups[Index]));
664}
665
666extern "C" void
667__llvm_profile_offload_register_section_shadow_variable(void *ptr) {
668 if (NumShadowVariables == 0)
669 return;
670
671 /* Match CGCUDANV.cpp: data, counters, uniform counters, then names for each
672 * kernel. */
673 OffloadSectionShadowGroup *Group =
674 &OffloadSectionShadowGroups[NumShadowVariables - 1];
675 int ShadowIndex = Group->NumSections / 4;
676 if (ensureSectionShadowCapacity(Group, MinCapacity: ShadowIndex + 1))
677 return;
678 if (ShadowIndex >= Group->NumShadows)
679 Group->NumShadows = ShadowIndex + 1;
680
681 OffloadSectionShadow *Shadow = &Group->Shadows[ShadowIndex];
682 switch (Group->NumSections % 4) {
683 case 0:
684 Shadow->Data = ptr;
685 break;
686 case 1:
687 Shadow->Counters = ptr;
688 break;
689 case 2:
690 Shadow->UniformCounters = ptr;
691 break;
692 case 3:
693 Shadow->Names = ptr;
694 break;
695 }
696 ++Group->NumSections;
697}
698
699namespace {
700
701struct ProfileSectionCopy {
702 const char *Name;
703 const void *DevBegin;
704 size_t Size;
705 const void *&CachedDevBegin;
706 char *&CachedHost;
707 size_t &CachedSize;
708 UniqueFree Owner;
709 char *HostBegin = nullptr;
710 bool Reused = false;
711
712 ProfileSectionCopy(const char *Name, const void *DevBegin, size_t Size,
713 const void *&CachedDevBegin, char *&CachedHost,
714 size_t &CachedSize)
715 : Name(Name), DevBegin(DevBegin), Size(Size),
716 CachedDevBegin(CachedDevBegin), CachedHost(CachedHost),
717 CachedSize(CachedSize) {}
718
719 ProfileSectionCopy(const ProfileSectionCopy &) = delete;
720 ProfileSectionCopy &operator=(const ProfileSectionCopy &) = delete;
721
722 int prepare() {
723 if (Size == 0)
724 return 0;
725 if (DevBegin == CachedDevBegin && Size == CachedSize) {
726 HostBegin = CachedHost;
727 Reused = true;
728 if (isVerboseMode())
729 PROF_NOTE("Reusing cached %s section (%zu bytes)\n", Name, Size);
730 } else {
731 HostBegin = static_cast<char *>(malloc(size: Size));
732 Owner.reset(P: HostBegin);
733 }
734 return HostBegin ? 0 : -1;
735 }
736
737 int copy() {
738 if (Size == 0 || Reused)
739 return 0;
740 return memcpyDeviceToHost(Dst: HostBegin, Src: DevBegin, Size);
741 }
742
743 void commitCache() {
744 if (Reused || Size == 0)
745 return;
746 CachedDevBegin = DevBegin;
747 CachedHost = HostBegin;
748 CachedSize = Size;
749 Owner.release();
750 }
751};
752
753} // namespace
754
755static int getRegisteredSectionBounds(void *Shadow, void **DevicePtr,
756 size_t *Size) {
757 *DevicePtr = nullptr;
758 *Size = 0;
759 int AddrRc = hipGetSymbolAddress(devPtr: DevicePtr, symbol: Shadow);
760 int SizeRc = hipGetSymbolSize(size: Size, symbol: Shadow);
761 return AddrRc == 0 && SizeRc == 0 && *DevicePtr && *Size > 0 ? 0 : -1;
762}
763
764struct RegisteredSectionRange {
765 const void *Data;
766 const void *Counters;
767 const void *UniformCounters;
768 const void *Names;
769 size_t DataSize;
770 size_t CountersSize;
771 size_t UniformCountersSize;
772 size_t NamesSize;
773 size_t DataOffset;
774 size_t CountersOffset;
775 size_t UniformCountersOffset;
776 size_t NamesOffset;
777};
778
779static int
780hasCompleteSectionShadows(const OffloadSectionShadowGroup *Sections) {
781 if (!Sections || Sections->NumShadows == 0 || Sections->NumSections % 4 != 0)
782 return 0;
783 for (int I = 0; I < Sections->NumShadows; ++I) {
784 if (!Sections->Shadows[I].Data || !Sections->Shadows[I].Counters ||
785 !Sections->Shadows[I].UniformCounters || !Sections->Shadows[I].Names)
786 return 0;
787 }
788 return 1;
789}
790
791int __prof_rocm::processDeviceOffloadPrf(
792 void *DeviceOffloadPrf, const char *Target,
793 const OffloadSectionShadowGroup *Sections) {
794 __llvm_profile_gpu_sections HostSections;
795
796 if (hipMemcpy(dest: &HostSections, src: DeviceOffloadPrf, len: sizeof(HostSections),
797 kind: 2 /*DToH*/) != 0) {
798 PROF_ERR("%s\n", "failed to copy offload prf structure from device");
799 return -1;
800 }
801
802 const void *DevCntsBegin = HostSections.CountersStart;
803 const void *DevDataBegin = HostSections.DataStart;
804 const void *DevNamesBegin = HostSections.NamesStart;
805 const void *DevUniformCntsBegin = HostSections.UniformCountersStart;
806 const void *DevCntsEnd = HostSections.CountersStop;
807 const void *DevDataEnd = HostSections.DataStop;
808 const void *DevNamesEnd = HostSections.NamesStop;
809 const void *DevUniformCntsEnd = HostSections.UniformCountersStop;
810
811 size_t CountersSize = (const char *)DevCntsEnd - (const char *)DevCntsBegin;
812 size_t DataSize = (const char *)DevDataEnd - (const char *)DevDataBegin;
813 size_t NamesSize = (const char *)DevNamesEnd - (const char *)DevNamesBegin;
814 size_t UniformCountersSize =
815 (const char *)DevUniformCntsEnd - (const char *)DevUniformCntsBegin;
816
817 int UseRegisteredSections = hasCompleteSectionShadows(Sections);
818 RegisteredSectionRange *RegisteredRanges = nullptr;
819 int NumRegisteredRanges = 0;
820
821 if (isVerboseMode())
822 PROF_NOTE("Section pointers: Cnts=[%p,%p]=%zu Data=[%p,%p]=%zu "
823 "Names=[%p,%p]=%zu UCnts=[%p,%p]=%zu\n",
824 DevCntsBegin, DevCntsEnd, CountersSize, DevDataBegin, DevDataEnd,
825 DataSize, DevNamesBegin, DevNamesEnd, NamesSize,
826 DevUniformCntsBegin, DevUniformCntsEnd, UniformCountersSize);
827
828 if (CountersSize == 0 || DataSize == 0)
829 return 0;
830
831 int ret = -1;
832
833 /* Sections using linker-defined __start_/__stop_ bounds are shared across
834 TU structs in RDC mode. Deduplicate by caching the last copied range. */
835 static const void *CachedDevNamesBegin = nullptr;
836 static char *CachedHostNames = nullptr;
837 static size_t CachedNamesSize = 0;
838
839 static const void *CachedDevCntsBegin = nullptr;
840 static char *CachedHostCnts = nullptr;
841 static size_t CachedCntsSize = 0;
842
843 static const void *CachedDevDataBegin = nullptr;
844 static char *CachedHostData = nullptr;
845 static size_t CachedDataSize = 0;
846
847 static const void *CachedDevUCntsBegin = nullptr;
848 static char *CachedHostUCnts = nullptr;
849 static size_t CachedUCntsSize = 0;
850
851 ProfileSectionCopy Cnts("counters", DevCntsBegin, CountersSize,
852 CachedDevCntsBegin, CachedHostCnts, CachedCntsSize);
853 ProfileSectionCopy Data("data", DevDataBegin, DataSize, CachedDevDataBegin,
854 CachedHostData, CachedDataSize);
855 ProfileSectionCopy Names("names", DevNamesBegin, NamesSize,
856 CachedDevNamesBegin, CachedHostNames,
857 CachedNamesSize);
858 ProfileSectionCopy UCnts("ucnts", DevUniformCntsBegin, UniformCountersSize,
859 CachedDevUCntsBegin, CachedHostUCnts,
860 CachedUCntsSize);
861
862 UniqueFree RegisteredRangeOwner;
863
864 if (UseRegisteredSections) {
865 NumRegisteredRanges = Sections->NumShadows;
866 RegisteredRangeOwner.reset(
867 P: malloc(size: NumRegisteredRanges * sizeof(RegisteredSectionRange)));
868 RegisteredRanges = (RegisteredSectionRange *)RegisteredRangeOwner.get();
869 if (!RegisteredRanges) {
870 PROF_ERR("%s\n", "failed to allocate registered section table");
871 return -1;
872 }
873 __builtin_memset(RegisteredRanges, 0,
874 NumRegisteredRanges * sizeof(*RegisteredRanges));
875
876 size_t RegisteredDataSize = 0;
877 size_t RegisteredCountersSize = 0;
878 size_t RegisteredUniformCountersSize = 0;
879 size_t RegisteredNamesSize = 0;
880 for (int I = 0; I < NumRegisteredRanges; ++I) {
881 void *Data = nullptr;
882 void *Counters = nullptr;
883 void *UniformCounters = nullptr;
884 void *Names = nullptr;
885 size_t ThisDataSize = 0;
886 size_t ThisCountersSize = 0;
887 size_t ThisUniformCountersSize = 0;
888 size_t ThisNamesSize = 0;
889 OffloadSectionShadow *Shadow = &Sections->Shadows[I];
890 if (getRegisteredSectionBounds(Shadow: Shadow->Data, DevicePtr: &Data, Size: &ThisDataSize) != 0 ||
891 getRegisteredSectionBounds(Shadow: Shadow->Counters, DevicePtr: &Counters,
892 Size: &ThisCountersSize) != 0 ||
893 getRegisteredSectionBounds(Shadow: Shadow->UniformCounters, DevicePtr: &UniformCounters,
894 Size: &ThisUniformCountersSize) != 0 ||
895 getRegisteredSectionBounds(Shadow: Shadow->Names, DevicePtr: &Names, Size: &ThisNamesSize) !=
896 0) {
897 PROF_ERR("%s\n", "failed to get registered section bounds");
898 return -1;
899 }
900
901 RegisteredRanges[I].Data = Data;
902 RegisteredRanges[I].Counters = Counters;
903 RegisteredRanges[I].UniformCounters = UniformCounters;
904 RegisteredRanges[I].Names = Names;
905 RegisteredRanges[I].DataSize = ThisDataSize;
906 RegisteredRanges[I].CountersSize = ThisCountersSize;
907 RegisteredRanges[I].UniformCountersSize = ThisUniformCountersSize;
908 RegisteredRanges[I].NamesSize = ThisNamesSize;
909 RegisteredRanges[I].DataOffset = RegisteredDataSize;
910 RegisteredRanges[I].CountersOffset = RegisteredCountersSize;
911 RegisteredRanges[I].UniformCountersOffset = RegisteredUniformCountersSize;
912 RegisteredDataSize += ThisDataSize;
913 RegisteredCountersSize += ThisCountersSize;
914 RegisteredUniformCountersSize += ThisUniformCountersSize;
915
916 int ReuseNames = 0;
917 for (int J = 0; J < I; ++J) {
918 if (RegisteredRanges[J].Names == Names &&
919 RegisteredRanges[J].NamesSize == ThisNamesSize) {
920 RegisteredRanges[I].NamesOffset = RegisteredRanges[J].NamesOffset;
921 ReuseNames = 1;
922 break;
923 }
924 }
925 if (!ReuseNames) {
926 RegisteredRanges[I].NamesOffset = RegisteredNamesSize;
927 RegisteredNamesSize += ThisNamesSize;
928 }
929 }
930
931 DataSize = RegisteredDataSize;
932 CountersSize = RegisteredCountersSize;
933 UniformCountersSize = RegisteredUniformCountersSize;
934 NamesSize = RegisteredNamesSize;
935 Data.HostBegin = DataSize ? (char *)malloc(size: DataSize) : nullptr;
936 Cnts.HostBegin = CountersSize ? (char *)malloc(size: CountersSize) : nullptr;
937 UCnts.HostBegin =
938 UniformCountersSize ? (char *)malloc(size: UniformCountersSize) : nullptr;
939 Names.HostBegin = NamesSize ? (char *)malloc(size: NamesSize) : nullptr;
940 Data.Owner.reset(P: Data.HostBegin);
941 Cnts.Owner.reset(P: Cnts.HostBegin);
942 UCnts.Owner.reset(P: UCnts.HostBegin);
943 Names.Owner.reset(P: Names.HostBegin);
944 if ((DataSize > 0 && !Data.HostBegin) ||
945 (CountersSize > 0 && !Cnts.HostBegin) ||
946 (UniformCountersSize > 0 && !UCnts.HostBegin) ||
947 (NamesSize > 0 && !Names.HostBegin)) {
948 PROF_ERR("%s\n", "failed to allocate host memory for device sections");
949 return -1;
950 }
951
952 for (int I = 0; I < NumRegisteredRanges; ++I) {
953 RegisteredSectionRange *R = &RegisteredRanges[I];
954 if (memcpyDeviceToHost(Dst: Data.HostBegin + R->DataOffset, Src: R->Data,
955 Size: R->DataSize) != 0 ||
956 memcpyDeviceToHost(Dst: Cnts.HostBegin + R->CountersOffset, Src: R->Counters,
957 Size: R->CountersSize) != 0 ||
958 memcpyDeviceToHost(Dst: UCnts.HostBegin + R->UniformCountersOffset,
959 Src: R->UniformCounters, Size: R->UniformCountersSize) != 0) {
960 PROF_ERR("%s\n", "failed to copy profile sections from device");
961 return -1;
962 }
963
964 int CopyNames = 1;
965 for (int J = 0; J < I; ++J) {
966 if (RegisteredRanges[J].Names == R->Names &&
967 RegisteredRanges[J].NamesSize == R->NamesSize) {
968 CopyNames = 0;
969 break;
970 }
971 }
972 if (CopyNames && R->NamesSize > 0 &&
973 memcpyDeviceToHost(Dst: Names.HostBegin + R->NamesOffset, Src: R->Names,
974 Size: R->NamesSize) != 0) {
975 PROF_ERR("%s\n", "failed to copy profile sections from device");
976 return -1;
977 }
978 }
979 } else {
980 if (Cnts.prepare() != 0 || Data.prepare() != 0 || Names.prepare() != 0 ||
981 UCnts.prepare() != 0) {
982 PROF_ERR("%s\n", "failed to allocate host memory for device sections");
983 return -1;
984 }
985
986 if (Data.copy() != 0 || Cnts.copy() != 0 || Names.copy() != 0 ||
987 UCnts.copy() != 0) {
988 PROF_ERR("%s\n", "failed to copy profile sections from device");
989 return -1;
990 }
991
992 /* Cache buffers so RDC-mode multi-shadow drains can reuse them.
993 * release() prevents the scope guards from freeing what the cache owns. */
994 Cnts.commitCache();
995 Data.commitCache();
996 Names.commitCache();
997 UCnts.commitCache();
998 }
999
1000 if (isVerboseMode())
1001 PROF_NOTE("Copied device sections: Counters=%zu, Data=%zu, Names=%zu, "
1002 "UniformCounters=%zu\n",
1003 CountersSize, DataSize, NamesSize, UniformCountersSize);
1004
1005 // Arrange buffer as [Data][Padding][Counters][Names] to match the layout
1006 // expected by lprofWriteDataImpl (CountersDelta = CountersBegin - DataBegin).
1007 const uint64_t NumData = DataSize / sizeof(__llvm_profile_data);
1008 const uint64_t NumBitmapBytes = 0;
1009 const uint64_t NumUniformCounters = UniformCountersSize / sizeof(uint64_t);
1010 const uint64_t VTableSectionSize = 0;
1011 const uint64_t VNamesSize = 0;
1012 uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
1013 PaddingBytesAfterBitmapBytes, PaddingBytesAfterUniformCounters,
1014 PaddingBytesAfterNames, PaddingBytesAfterVTable, PaddingBytesAfterVNames;
1015
1016 if (__llvm_profile_get_padding_sizes_for_counters(
1017 DataSize, CountersSize, NumBitmapBytes, NumUniformCounters, NamesSize,
1018 VTableSize: VTableSectionSize, VNameSize: VNamesSize, PaddingBytesBeforeCounters: &PaddingBytesBeforeCounters,
1019 PaddingBytesAfterCounters: &PaddingBytesAfterCounters, PaddingBytesAfterBitmap: &PaddingBytesAfterBitmapBytes,
1020 PaddingBytesAfterUniformCounters: &PaddingBytesAfterUniformCounters, PaddingBytesAfterNames: &PaddingBytesAfterNames,
1021 PaddingBytesAfterVTable: &PaddingBytesAfterVTable, PaddingBytesAfterVNames: &PaddingBytesAfterVNames) != 0) {
1022 PROF_ERR("%s\n", "failed to get padding sizes");
1023 return -1;
1024 }
1025
1026 size_t ContiguousBufferSize =
1027 DataSize + PaddingBytesBeforeCounters + CountersSize + NamesSize;
1028 UniqueFree ContiguousBuf(malloc(size: ContiguousBufferSize));
1029 if (!ContiguousBuf.get()) {
1030 PROF_ERR("%s\n", "failed to allocate contiguous buffer");
1031 return -1;
1032 }
1033 char *ContiguousBuffer = ContiguousBuf.get();
1034 __builtin_memset(ContiguousBuffer, 0, ContiguousBufferSize);
1035
1036 char *BufDataBegin = ContiguousBuffer;
1037 char *BufCountersBegin =
1038 ContiguousBuffer + DataSize + PaddingBytesBeforeCounters;
1039 char *BufNamesBegin = BufCountersBegin + CountersSize;
1040
1041 __builtin_memcpy(BufDataBegin, Data.HostBegin, DataSize);
1042 __builtin_memcpy(BufCountersBegin, Cnts.HostBegin, CountersSize);
1043 __builtin_memcpy(BufNamesBegin, Names.HostBegin, NamesSize);
1044
1045 // CounterPtr and UniformCounterPtr are device-relative offsets; relocate
1046 // them for the file layout where the Data section precedes the Counters and
1047 // UniformCounters sections. Uniform counters are copied in linker (section)
1048 // order and located via their relative pointer, exactly like the regular
1049 // counters: llvm-profdata reads them through UniformCounterPtr (decrementing
1050 // UniformCountersDelta per record, just like CountersDelta) and does not
1051 // assume data-record order, so no reordering is needed.
1052 ptrdiff_t UCFileOffset = DataSize + PaddingBytesBeforeCounters +
1053 CountersSize + PaddingBytesAfterCounters +
1054 NumBitmapBytes + PaddingBytesAfterBitmapBytes;
1055 __llvm_profile_data *RelocatedData = (__llvm_profile_data *)BufDataBegin;
1056 for (uint64_t i = 0; i < NumData; ++i) {
1057 size_t DataRecordOffset = i * sizeof(__llvm_profile_data);
1058 const char *RangeDevDataBegin = (const char *)DevDataBegin;
1059 const char *RangeDevCountersBegin = (const char *)DevCntsBegin;
1060 const char *RangeDevUCntsBegin = (const char *)DevUniformCntsBegin;
1061 size_t RangeCountersOffset = 0;
1062 size_t RangeUCntsOffset = 0;
1063 if (UseRegisteredSections) {
1064 int FoundRange = 0;
1065 for (int R = 0; R < NumRegisteredRanges; ++R) {
1066 RegisteredSectionRange *Range = &RegisteredRanges[R];
1067 if (DataRecordOffset < Range->DataOffset ||
1068 DataRecordOffset >= Range->DataOffset + Range->DataSize)
1069 continue;
1070 RangeDevDataBegin = (const char *)Range->Data;
1071 RangeDevCountersBegin = (const char *)Range->Counters;
1072 RangeDevUCntsBegin = (const char *)Range->UniformCounters;
1073 RangeCountersOffset = Range->CountersOffset;
1074 RangeUCntsOffset = Range->UniformCountersOffset;
1075 DataRecordOffset -= Range->DataOffset;
1076 FoundRange = 1;
1077 break;
1078 }
1079 if (!FoundRange) {
1080 PROF_ERR("%s\n", "failed to locate profile data record range");
1081 return -1;
1082 }
1083 }
1084 const char *DeviceDataStructAddr = RangeDevDataBegin + DataRecordOffset;
1085 if (RelocatedData[i].CounterPtr) {
1086 const char *DeviceCountersAddr =
1087 DeviceDataStructAddr + (ptrdiff_t)RelocatedData[i].CounterPtr;
1088 ptrdiff_t OffsetIntoCountersSection =
1089 DeviceCountersAddr - RangeDevCountersBegin;
1090 ptrdiff_t NewRelativeOffset =
1091 DataSize + PaddingBytesBeforeCounters + RangeCountersOffset +
1092 OffsetIntoCountersSection - (i * sizeof(__llvm_profile_data));
1093 __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
1094 offsetof(__llvm_profile_data, CounterPtr),
1095 &NewRelativeOffset, sizeof(NewRelativeOffset));
1096 }
1097 if (UCnts.HostBegin && RelocatedData[i].UniformCounterPtr) {
1098 const char *DeviceUCAddr =
1099 DeviceDataStructAddr + (ptrdiff_t)RelocatedData[i].UniformCounterPtr;
1100 ptrdiff_t OffsetIntoUCSection = DeviceUCAddr - RangeDevUCntsBegin;
1101 ptrdiff_t NewUCRelativeOffset = UCFileOffset + RangeUCntsOffset +
1102 OffsetIntoUCSection -
1103 (i * sizeof(__llvm_profile_data));
1104 __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
1105 offsetof(__llvm_profile_data, UniformCounterPtr),
1106 &NewUCRelativeOffset, sizeof(NewUCRelativeOffset));
1107 } else {
1108 __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
1109 offsetof(__llvm_profile_data, UniformCounterPtr),
1110 0, sizeof(RelocatedData[i].UniformCounterPtr));
1111 }
1112 __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
1113 offsetof(__llvm_profile_data, BitmapPtr),
1114 0,
1115 sizeof(RelocatedData[i].BitmapPtr) +
1116 sizeof(RelocatedData[i].FunctionPointer) +
1117 sizeof(RelocatedData[i].Values));
1118 }
1119
1120 ret = __llvm_write_custom_profile(
1121 Target, DataBegin: (__llvm_profile_data *)BufDataBegin,
1122 DataEnd: (__llvm_profile_data *)(BufDataBegin + DataSize), CountersBegin: BufCountersBegin,
1123 CountersEnd: BufCountersBegin + CountersSize, UniformCountersBegin: UCnts.HostBegin,
1124 UniformCountersEnd: UCnts.HostBegin ? UCnts.HostBegin + UniformCountersSize : nullptr,
1125 NamesBegin: BufNamesBegin, NamesEnd: BufNamesBegin + NamesSize, VersionOverride: nullptr);
1126
1127 if (ret != 0) {
1128 PROF_ERR("%s\n", "failed to write device profile using shared API");
1129 } else {
1130#if defined(__linux__) && !defined(_WIN32)
1131 // Dedup against the supplemental HSA pass: this section is now drained, so
1132 // the HSA walk must not drain the same device code object again.
1133 profRecordDrainedBounds(Data: DevDataBegin, Counters: DevCntsBegin, Names: DevNamesBegin);
1134#endif
1135 if (isVerboseMode())
1136 PROF_NOTE("%s\n", "Successfully wrote device profile using shared API");
1137 }
1138
1139 return ret;
1140}
1141
1142static int processShadowVariable(int Index, const char *Target) {
1143 void *ShadowVar = OffloadShadowVariables[Index];
1144 void *DeviceSections = nullptr;
1145 if (hipGetSymbolAddress(devPtr: &DeviceSections, symbol: ShadowVar) != 0) {
1146 PROF_WARN("failed to get symbol address for shadow variable %p\n",
1147 ShadowVar);
1148 return -1;
1149 }
1150 /* DeviceSections points at the per-TU sections struct itself. */
1151 const OffloadSectionShadowGroup *Sections = nullptr;
1152 if (Index < CapSectionShadowGroups)
1153 Sections = &OffloadSectionShadowGroups[Index];
1154 if (!hasCompleteSectionShadows(Sections))
1155 return 0;
1156 return processDeviceOffloadPrf(DeviceOffloadPrf: DeviceSections, Target, Sections);
1157}
1158
1159static int isHipAvailable(void) {
1160 ensureHipLoaded();
1161 return pHipMemcpy != nullptr && pHipGetSymbolAddress != nullptr;
1162}
1163
1164/* -------------------------------------------------------------------------- */
1165/* Collect device-side profile data */
1166/* -------------------------------------------------------------------------- */
1167
1168/* Host-shadow drain: static-linked kernels (host __hipRegisterVar shadows) and
1169 * intercepted dynamic modules. The caller gates this on
1170 * (NumShadowVariables || NumDynamicModules) && isHipAvailable(); pure
1171 * device-linked programs (RCCL) are handled by the supplemental HSA pass. */
1172static int collectHostShadowData(void) {
1173 int Ret = 0;
1174
1175 /* Shadow variables (static-linked kernels): drain from every device. */
1176 if (NumShadowVariables > 0) {
1177 int OrigDevice = -1;
1178 hipGetDevice(DeviceId: &OrigDevice);
1179
1180 for (int Dev = 0; Dev < NumDevices; ++Dev) {
1181 if (!shouldCollectDevice(DeviceId: Dev)) {
1182 if (isVerboseMode())
1183 PROF_NOTE("Skipping unused device %d\n", Dev);
1184 continue;
1185 }
1186#if defined(__linux__) && !defined(_WIN32)
1187 /* When no kernel launch was tracked at all, shouldCollectDevice() falls
1188 * back to collect-all, which can fault/hang reading a non-resident
1189 * device's sections on a multi-GPU host. On Linux the supplemental HSA
1190 * drain covers those cases safely. */
1191 if (!__atomic_load_n(&AnyDeviceUsed, __ATOMIC_ACQUIRE)) {
1192 if (isVerboseMode())
1193 PROF_NOTE("No tracked launch; deferring device %d to HSA drain\n",
1194 Dev);
1195 continue;
1196 }
1197#endif
1198 if (hipSetDevice(DeviceId: Dev) != 0) {
1199 if (isVerboseMode())
1200 PROF_NOTE("Failed to set device %d, skipping\n", Dev);
1201 continue;
1202 }
1203 const char *ArchName = getDeviceArchName(DeviceId: Dev);
1204 if (isVerboseMode())
1205 PROF_NOTE("Collecting static profile data from device %d (%s)\n", Dev,
1206 ArchName);
1207 for (int i = 0; i < NumShadowVariables; ++i) {
1208 /* Stable name per shadow so a repeated drain (explicit collect plus the
1209 * atexit drain) overwrites its own profraw rather than emitting a
1210 * second one: bare arch for a single TU, arch.<i> for RDC multi-TU. */
1211 const char *Target = ArchName;
1212 char TargetWithIdx[64];
1213 if (NumShadowVariables > 1) {
1214 snprintf(s: TargetWithIdx, maxlen: sizeof(TargetWithIdx), format: "%s.%d", ArchName, i);
1215 Target = TargetWithIdx;
1216 }
1217 if (processShadowVariable(Index: i, Target) != 0)
1218 Ret = -1;
1219 }
1220 }
1221
1222 if (OrigDevice >= 0)
1223 hipSetDevice(DeviceId: OrigDevice);
1224 }
1225
1226 /* Warn about unprocessed TUs; skip cleared slots (already drained). */
1227 lockDynamicModules();
1228 for (int i = 0; i < NumDynamicModules; ++i) {
1229 OffloadDynamicModuleInfo *MI = &DynamicModules[i];
1230 if (!MI->ModulePtr)
1231 continue;
1232 for (int t = 0; t < MI->NumTUs; ++t) {
1233 if (!MI->TUs[t].Processed) {
1234 PROF_WARN("dynamic module %p TU %d was not processed before exit\n",
1235 MI->ModulePtr, t);
1236 Ret = -1;
1237 }
1238 }
1239 }
1240 unlockDynamicModules();
1241
1242 return Ret;
1243}
1244
1245extern "C" int __llvm_profile_hip_collect_device_data(void) {
1246 int Ret = 0;
1247
1248 if ((NumShadowVariables != 0 || NumDynamicModules != 0) && isHipAvailable() &&
1249 collectHostShadowData() != 0)
1250 Ret = -1;
1251
1252#if defined(__linux__) && !defined(_WIN32)
1253 /* Supplemental HSA-introspection drain */
1254 if (drainDevicesViaHsa() != 0)
1255 Ret = -1;
1256#endif
1257
1258 if (Ret != 0)
1259 PROF_WARN("%s\n", "failed to collect device profile data");
1260 return Ret;
1261}
1262
1263/* Linux HIP interceptors. */
1264
1265#if defined(__linux__) && !defined(_WIN32)
1266
1267typedef struct {
1268 unsigned int x;
1269 unsigned int y;
1270 unsigned int z;
1271} HipDim3;
1272
1273typedef struct {
1274 void *Func;
1275 HipDim3 GridDim;
1276 HipDim3 BlockDim;
1277 void **Args;
1278 size_t SharedMem;
1279 HipStream Stream;
1280} HipLaunchParams;
1281
1282typedef struct {
1283 HipDim3 GridDim;
1284 HipDim3 BlockDim;
1285 size_t DynamicSmemBytes;
1286 HipStream Stream;
1287 void *Attrs;
1288 unsigned NumAttrs;
1289} HipLaunchConfig;
1290
1291typedef void *HipFunction;
1292typedef void *HipEvent;
1293typedef void *HipGraphExec;
1294
1295static int recordHipLaunchResult(int Rc, HipStream Stream) {
1296 if (Rc == 0)
1297 markLaunchStreamDeviceUsed(Stream);
1298 return Rc;
1299}
1300
1301static int recordHipMultiDeviceLaunchResult(int Rc,
1302 HipLaunchParams *LaunchParams,
1303 int NumLaunches) {
1304 if (Rc != 0 || !LaunchParams || NumLaunches <= 0)
1305 return Rc;
1306 for (int I = 0; I < NumLaunches; ++I)
1307 markLaunchStreamDeviceUsed(Stream: LaunchParams[I].Stream);
1308 return Rc;
1309}
1310
1311// interceptors must have external linkage
1312// NOLINTBEGIN(misc-use-internal-linkage)
1313INTERCEPTOR(int, hipLaunchKernel, const void *Function, HipDim3 GridDim,
1314 HipDim3 BlockDim, void **Args, size_t SharedMemBytes,
1315 HipStream Stream) {
1316 return recordHipLaunchResult(REAL(hipLaunchKernel)(Function, GridDim,
1317 BlockDim, Args,
1318 SharedMemBytes, Stream),
1319 Stream);
1320}
1321
1322INTERCEPTOR(int, hipLaunchKernel_spt, const void *Function, HipDim3 GridDim,
1323 HipDim3 BlockDim, void **Args, size_t SharedMemBytes,
1324 HipStream Stream) {
1325 return recordHipLaunchResult(
1326 REAL(hipLaunchKernel_spt)(Function, GridDim, BlockDim, Args,
1327 SharedMemBytes, Stream),
1328 Stream);
1329}
1330
1331INTERCEPTOR(int, hipExtLaunchKernel, const void *Function, HipDim3 GridDim,
1332 HipDim3 BlockDim, void **Args, size_t SharedMemBytes,
1333 HipStream Stream, HipEvent StartEvent, HipEvent StopEvent,
1334 int Flags) {
1335 return recordHipLaunchResult(
1336 REAL(hipExtLaunchKernel)(Function, GridDim, BlockDim, Args,
1337 SharedMemBytes, Stream, StartEvent, StopEvent,
1338 Flags),
1339 Stream);
1340}
1341
1342INTERCEPTOR(int, hipLaunchKernelExC, const HipLaunchConfig *Config,
1343 const void *Function, void **Args) {
1344 int Rc = REAL(hipLaunchKernelExC)(Config, Function, Args);
1345 return recordHipLaunchResult(Rc, Stream: Config ? Config->Stream : nullptr);
1346}
1347
1348INTERCEPTOR(int, hipLaunchCooperativeKernel, const void *Function,
1349 HipDim3 GridDim, HipDim3 BlockDim, void **KernelParams,
1350 unsigned SharedMemBytes, HipStream Stream) {
1351 return recordHipLaunchResult(
1352 REAL(hipLaunchCooperativeKernel)(Function, GridDim, BlockDim,
1353 KernelParams, SharedMemBytes, Stream),
1354 Stream);
1355}
1356
1357INTERCEPTOR(int, hipLaunchCooperativeKernel_spt, const void *Function,
1358 HipDim3 GridDim, HipDim3 BlockDim, void **KernelParams,
1359 unsigned SharedMemBytes, HipStream Stream) {
1360 return recordHipLaunchResult(
1361 REAL(hipLaunchCooperativeKernel_spt)(
1362 Function, GridDim, BlockDim, KernelParams, SharedMemBytes, Stream),
1363 Stream);
1364}
1365
1366INTERCEPTOR(int, hipLaunchCooperativeKernelMultiDevice,
1367 HipLaunchParams *LaunchParams, int NumDevices, unsigned Flags) {
1368 return recordHipMultiDeviceLaunchResult(
1369 REAL(hipLaunchCooperativeKernelMultiDevice)(LaunchParams, NumDevices,
1370 Flags),
1371 LaunchParams, NumLaunches: NumDevices);
1372}
1373
1374INTERCEPTOR(int, hipExtLaunchMultiKernelMultiDevice,
1375 HipLaunchParams *LaunchParams, int NumDevices, unsigned Flags) {
1376 return recordHipMultiDeviceLaunchResult(
1377 REAL(hipExtLaunchMultiKernelMultiDevice)(LaunchParams, NumDevices, Flags),
1378 LaunchParams, NumLaunches: NumDevices);
1379}
1380
1381INTERCEPTOR(int, hipModuleLaunchKernel, HipFunction Function, unsigned GridDimX,
1382 unsigned GridDimY, unsigned GridDimZ, unsigned BlockDimX,
1383 unsigned BlockDimY, unsigned BlockDimZ, unsigned SharedMemBytes,
1384 HipStream Stream, void **KernelParams, void **Extra) {
1385 return recordHipLaunchResult(
1386 REAL(hipModuleLaunchKernel)(Function, GridDimX, GridDimY, GridDimZ,
1387 BlockDimX, BlockDimY, BlockDimZ,
1388 SharedMemBytes, Stream, KernelParams, Extra),
1389 Stream);
1390}
1391
1392INTERCEPTOR(int, hipExtModuleLaunchKernel, HipFunction Function,
1393 unsigned GridDimX, unsigned GridDimY, unsigned GridDimZ,
1394 unsigned BlockDimX, unsigned BlockDimY, unsigned BlockDimZ,
1395 size_t SharedMemBytes, HipStream Stream, void **KernelParams,
1396 void **Extra, HipEvent StartEvent, HipEvent StopEvent,
1397 unsigned Flags) {
1398 return recordHipLaunchResult(
1399 REAL(hipExtModuleLaunchKernel)(Function, GridDimX, GridDimY, GridDimZ,
1400 BlockDimX, BlockDimY, BlockDimZ,
1401 SharedMemBytes, Stream, KernelParams,
1402 Extra, StartEvent, StopEvent, Flags),
1403 Stream);
1404}
1405
1406INTERCEPTOR(int, hipGraphLaunch, HipGraphExec GraphExec, HipStream Stream) {
1407 return recordHipLaunchResult(REAL(hipGraphLaunch)(GraphExec, Stream), Stream);
1408}
1409
1410INTERCEPTOR(int, hipGraphLaunch_spt, HipGraphExec GraphExec, HipStream Stream) {
1411 return recordHipLaunchResult(REAL(hipGraphLaunch_spt)(GraphExec, Stream),
1412 Stream);
1413}
1414
1415static void registerDynamicModuleFromFile(int Rc, void **Module,
1416 const char *Filename) {
1417 if (Rc)
1418 return;
1419
1420 FILE *File = fopen(filename: Filename, modes: "rb");
1421 if (!File) {
1422 PROF_WARN("failed to open dynamic module file %s\n", Filename);
1423 return;
1424 }
1425
1426 if (fseek(stream: File, off: 0, SEEK_END) != 0) {
1427 PROF_WARN("failed to seek dynamic module file %s\n", Filename);
1428 fclose(stream: File);
1429 return;
1430 }
1431 long FileSize = ftell(stream: File);
1432 if (FileSize <= 0 || fseek(stream: File, off: 0, SEEK_SET) != 0) {
1433 PROF_WARN("failed to get size of dynamic module file %s\n", Filename);
1434 fclose(stream: File);
1435 return;
1436 }
1437
1438 ManagedMemory Image;
1439 lprofGetFileContentBuffer(F: File, FileSize: (uint64_t)FileSize, Buf: &Image);
1440 fclose(stream: File);
1441 if (Image.Status == MS_INVALID) {
1442 PROF_WARN("failed to read dynamic module file %s\n", Filename);
1443 return;
1444 }
1445
1446 __llvm_profile_offload_register_dynamic_module(ModuleLoadRc: Rc, Ptr: Module, Image: Image.Addr);
1447 lprofReleaseBuffer(FileBuffer: &Image, Length: (size_t)FileSize);
1448}
1449
1450INTERCEPTOR(int, hipModuleLoad, void **module, const char *fname) {
1451 int rc = REAL(hipModuleLoad)(module, fname);
1452 registerDynamicModuleFromFile(Rc: rc, Module: module, Filename: fname);
1453 return rc;
1454}
1455
1456INTERCEPTOR(int, hipModuleLoadData, void **module, const void *image) {
1457 int rc = REAL(hipModuleLoadData)(module, image);
1458 __llvm_profile_offload_register_dynamic_module(ModuleLoadRc: rc, Ptr: module, Image: image);
1459 return rc;
1460}
1461
1462INTERCEPTOR(int, hipModuleLoadDataEx, void **module, const void *image,
1463 unsigned numOptions, void **options, void **optionValues) {
1464 int rc = REAL(hipModuleLoadDataEx)(module, image, numOptions, options,
1465 optionValues);
1466 __llvm_profile_offload_register_dynamic_module(ModuleLoadRc: rc, Ptr: module, Image: image);
1467 return rc;
1468}
1469
1470INTERCEPTOR(int, hipModuleUnload, void *module) {
1471 /* Drain counters before the module is destroyed; device addresses
1472 * captured at register time are invalid after unload. */
1473 __llvm_profile_offload_unregister_dynamic_module(Ptr: module);
1474 return REAL(hipModuleUnload)(module);
1475}
1476// NOLINTEND(misc-use-internal-linkage)
1477
1478__attribute__((constructor)) static void installHipInterceptors() {
1479 /* Avoid interception unless the HIP runtime is already loaded. */
1480 int HasModuleLoad = dlsym(RTLD_DEFAULT, name: "hipModuleLoad") != nullptr;
1481 int InstalledLaunch = 0;
1482#define TRY_INTERCEPT_LAUNCH(Name) \
1483 do { \
1484 if (dlsym(RTLD_DEFAULT, #Name)) \
1485 InstalledLaunch |= INTERCEPT_FUNCTION(Name); \
1486 } while (0)
1487 TRY_INTERCEPT_LAUNCH(hipLaunchKernel);
1488 TRY_INTERCEPT_LAUNCH(hipLaunchKernel_spt);
1489 TRY_INTERCEPT_LAUNCH(hipExtLaunchKernel);
1490 TRY_INTERCEPT_LAUNCH(hipLaunchKernelExC);
1491 TRY_INTERCEPT_LAUNCH(hipLaunchCooperativeKernel);
1492 TRY_INTERCEPT_LAUNCH(hipLaunchCooperativeKernel_spt);
1493 TRY_INTERCEPT_LAUNCH(hipLaunchCooperativeKernelMultiDevice);
1494 TRY_INTERCEPT_LAUNCH(hipExtLaunchMultiKernelMultiDevice);
1495 TRY_INTERCEPT_LAUNCH(hipModuleLaunchKernel);
1496 TRY_INTERCEPT_LAUNCH(hipExtModuleLaunchKernel);
1497 TRY_INTERCEPT_LAUNCH(hipGraphLaunch);
1498 TRY_INTERCEPT_LAUNCH(hipGraphLaunch_spt);
1499#undef TRY_INTERCEPT_LAUNCH
1500 int InstalledAny = InstalledLaunch;
1501 if (HasModuleLoad) {
1502 HasModuleLoad = INTERCEPT_FUNCTION(hipModuleLoad);
1503 InstalledAny |= HasModuleLoad;
1504 }
1505 if (!InstalledAny)
1506 return;
1507 if (isVerboseMode())
1508 PROF_NOTE("%s", "Installing HIP interceptors\n");
1509 if (HasModuleLoad) {
1510 INTERCEPT_FUNCTION(hipModuleLoadData);
1511 INTERCEPT_FUNCTION(hipModuleLoadDataEx);
1512 INTERCEPT_FUNCTION(hipModuleUnload);
1513 }
1514}
1515
1516#endif /* __linux__ */
1517