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