Reland HIP offload PGO runtime support as a separate opt-in library - #201606
Merged
yxsamliu merged 2 commits intoJun 5, 2026
Conversation
Relands the compiler-rt portion of llvm#177665 ("[PGO][AMDGPU] Add basic HIP offload PGO support"), which was reverted in llvm#201416. This restores the host-side ROCm/HIP device profile collection runtime (InstrProfilingPlatformROCm.cpp) and its build wiring as merged, including the follow-ups llvm#200101, llvm#200127, llvm#200111, and llvm#200859. This commit restores the reverted state verbatim; the next commit refactors the ROCm support into a separate opt-in library so the base profile library is not affected.
Contributor
Author
|
This is the compiler-rt half of the offload PGO reland. The compiler-side change (instrumentation, codegen, and the driver wiring that links clang_rt.profile_rocm) is in #201607 and depends on this one. |
|
@llvm/pr-subscribers-pgo Author: Yaxun (Sam) Liu (yxsamliu) ChangesThis relands the compiler-rt portion of #177665 ("[PGO][AMDGPU] Add basic HIP offload PGO support"), which was reverted in #201416, and then restructures the ROCm support so it no longer affects the base profile library. The first commit restores the reverted state verbatim, including the merged follow-ups #200101, #200127, #200111, and #200859. This brings back the host-side ROCm/HIP device profile collection runtime (InstrProfilingPlatformROCm.cpp) and its build wiring. The second commit moves the ROCm support out of clang_rt.profile into a separate, opt-in library clang_rt.profile_rocm. The ROCm collection runtime depends on the sanitizer interception and sanitizer_common object libraries, which are built with the dynamic CRT (/MD) on Windows. Merging them into clang_rt.profile forced the whole profile archive to /MD and broke existing static-CRT (/MT) consumers, which is why #177665 was reverted. Keeping ROCm support in a dedicated archive lets it be built the way ROCm needs (/MD on Windows, with the interceptor dependencies) while clang_rt.profile stays exactly as before: /MT on Windows, no ROCm dependency, no interceptor merge. The base runtime gains one optional hook: a function pointer lprofCollectDeviceProfilingData (null by default, the same pattern as FreeHook), which __llvm_profile_write_file calls after writing the host profile. clang_rt.profile_rocm registers its collector through this hook from a load-time constructor, so the base library has no reference to any ROCm symbol and the previous weak/strong-extern handling is removed. clang_rt.profile_rocm is gated on COMPILER_RT_BUILD_PROFILE_ROCM, which defaults off on all platforms, so a default build is never affected. The driver wiring to link it for HIP -fprofile-generate, and the device-side instrumentation that produces the profile data it collects, come in a separate change. Patch is 38.70 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/201606.diff 6 Files Affected:
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index e88321d822f84..fdcf3a64f73f9 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -322,6 +322,16 @@ option(COMPILER_RT_USE_ATOMIC_LIBRARY "Use compiler-rt atomic instead of libatom
option(COMPILER_RT_PROFILE_BAREMETAL "Build minimal baremetal profile library" OFF)
+# Opt-in: build clang_rt.profile_rocm, the host-side ROCm/HIP device profile
+# collection runtime. Kept separate from clang_rt.profile because it depends on
+# the sanitizer interception library (built /MD on Windows) and must not change
+# how the base profile library is built. Off by default on all platforms, so a
+# default build is never affected; users targeting HIP device PGO opt in.
+option(COMPILER_RT_BUILD_PROFILE_ROCM
+ "Build the host-side ROCm/HIP device profile collection runtime (clang_rt.profile_rocm)"
+ OFF)
+mark_as_advanced(COMPILER_RT_BUILD_PROFILE_ROCM)
+
include(config-ix)
#================================
diff --git a/compiler-rt/lib/profile/CMakeLists.txt b/compiler-rt/lib/profile/CMakeLists.txt
index 8d9a773412a22..87e166cd1e8fd 100644
--- a/compiler-rt/lib/profile/CMakeLists.txt
+++ b/compiler-rt/lib/profile/CMakeLists.txt
@@ -215,3 +215,49 @@ else()
LINK_LIBS ${PROFILE_LINK_LIBS}
PARENT_TARGET profile)
endif()
+
+# clang_rt.profile_rocm: opt-in host-side ROCm/HIP device profile collection
+# runtime. Kept as a separate archive from clang_rt.profile because it pulls in
+# the sanitizer interception + sanitizer_common object libs (built /MD on
+# Windows). The driver links it on the host link only for HIP -fprofile-generate;
+# at load time its constructor registers the collector with the base runtime via
+# lprofCollectDeviceProfilingData, so the base library has no ROCm dependency.
+if(COMPILER_RT_BUILD_PROFILE_ROCM AND NOT COMPILER_RT_PROFILE_BAREMETAL)
+ set(PROFILE_ROCM_SOURCES InstrProfilingPlatformROCm.cpp)
+ set(PROFILE_ROCM_FLAGS ${EXTRA_FLAGS})
+ # C++ source: no exception personality (host archive is linked from C).
+ append_list_if(COMPILER_RT_HAS_FNO_EXCEPTIONS_FLAG -fno-exceptions
+ PROFILE_ROCM_FLAGS)
+
+ # The interceptor path references __sanitizer_internal_{memcpy,memset,...} and
+ # other sanitizer_common symbols; merge the same object libs as clang_rt.cfi
+ # (without coverage/symbolizer) so links stay self-contained. These targets
+ # only exist when sanitizers/memprof/xray/ctx-profile are enabled.
+ set(PROFILE_ROCM_OBJECT_LIBS)
+ if(COMPILER_RT_HAS_INTERCEPTION
+ AND TARGET RTInterception.${COMPILER_RT_DEFAULT_TARGET_ARCH}
+ AND TARGET RTSanitizerCommon.${COMPILER_RT_DEFAULT_TARGET_ARCH}
+ AND TARGET RTSanitizerCommonLibc.${COMPILER_RT_DEFAULT_TARGET_ARCH})
+ list(APPEND PROFILE_ROCM_OBJECT_LIBS
+ RTInterception
+ RTSanitizerCommon
+ RTSanitizerCommonLibc)
+ endif()
+
+ if(MSVC)
+ # The merged interception / sanitizer_common object libs are built with the
+ # dynamic CRT (/MD); match it here so this archive is self-consistent. This
+ # does not affect clang_rt.profile, which stays on /MT.
+ set(CMAKE_MSVC_RUNTIME_LIBRARY MultiThreadedDLL)
+ endif()
+
+ add_compiler_rt_runtime(clang_rt.profile_rocm
+ STATIC
+ ARCHS ${PROFILE_SUPPORTED_ARCH}
+ OBJECT_LIBS ${PROFILE_ROCM_OBJECT_LIBS}
+ CFLAGS ${PROFILE_ROCM_FLAGS}
+ SOURCES ${PROFILE_ROCM_SOURCES}
+ ADDITIONAL_HEADERS ${PROFILE_HEADERS}
+ LINK_LIBS ${PROFILE_LINK_LIBS}
+ PARENT_TARGET profile)
+endif()
diff --git a/compiler-rt/lib/profile/InstrProfilingFile.c b/compiler-rt/lib/profile/InstrProfilingFile.c
index 71127b05aafb8..21b671b49b737 100644
--- a/compiler-rt/lib/profile/InstrProfilingFile.c
+++ b/compiler-rt/lib/profile/InstrProfilingFile.c
@@ -1198,6 +1198,11 @@ int __llvm_profile_write_file(void) {
if (rc)
PROF_ERR("Failed to write file \"%s\": %s\n", Filename, strerror(errno));
+ /* Collect device-side profile data if a device-profile runtime registered a
+ * collector (e.g. clang_rt.profile_rocm for HIP). No-op otherwise. */
+ if (lprofCollectDeviceProfilingData)
+ (void)lprofCollectDeviceProfilingData();
+
// Restore SIGKILL.
if (PDeathSig == 1)
lprofRestoreSigKill();
diff --git a/compiler-rt/lib/profile/InstrProfilingInternal.h b/compiler-rt/lib/profile/InstrProfilingInternal.h
index 5647782527eb7..5f8ccba29abb8 100644
--- a/compiler-rt/lib/profile/InstrProfilingInternal.h
+++ b/compiler-rt/lib/profile/InstrProfilingInternal.h
@@ -190,6 +190,12 @@ uint64_t lprofGetLoadModuleSignature(void);
unsigned lprofProfileDumped(void);
void lprofSetProfileDumped(unsigned);
+/* Optional hook for collecting device-side profile data (e.g. HIP/ROCm).
+ * Null in the base runtime; an opt-in device-profile runtime
+ * (clang_rt.profile_rocm) registers its collector at load time. When set,
+ * __llvm_profile_write_file invokes it after writing the host profile. */
+COMPILER_RT_VISIBILITY extern int (*lprofCollectDeviceProfilingData)(void);
+
COMPILER_RT_VISIBILITY extern void (*FreeHook)(void *);
COMPILER_RT_VISIBILITY extern uint8_t *DynamicBufferIOBuffer;
COMPILER_RT_VISIBILITY extern uint32_t VPBufferSize;
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
new file mode 100644
index 0000000000000..3c14c386cba9a
--- /dev/null
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -0,0 +1,904 @@
+//===- InstrProfilingPlatformROCm.cpp - Profile data ROCm platform -------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+extern "C" {
+#include "InstrProfiling.h"
+#include "InstrProfilingInternal.h"
+#include "InstrProfilingPort.h"
+}
+
+#include "interception/interception.h"
+// C library headers (not <cstdio> etc.): clang_rt.profile is built with
+// -nostdinc++ and avoids the C++ standard library (see profile/CMakeLists.txt).
+#include <stddef.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#ifdef _WIN32
+#define WIN32_LEAN_AND_MEAN
+#include <windows.h>
+#else
+#include <dlfcn.h>
+#include <pthread.h>
+#endif
+
+/* Serialize one-time HIP loader resolution and DynamicModules mutations.
+ * Inline to avoid a sanitizer_common dependency. */
+#ifdef _WIN32
+static INIT_ONCE HipLoadedOnce = INIT_ONCE_STATIC_INIT;
+static CRITICAL_SECTION DynamicModulesLock;
+static INIT_ONCE DynamicModulesLockInit = INIT_ONCE_STATIC_INIT;
+static BOOL CALLBACK initDynamicModulesLockCb(PINIT_ONCE, PVOID, PVOID *) {
+ InitializeCriticalSection(&DynamicModulesLock);
+ return TRUE;
+}
+static void lockDynamicModules(void) {
+ InitOnceExecuteOnce(&DynamicModulesLockInit, initDynamicModulesLockCb, NULL,
+ NULL);
+ EnterCriticalSection(&DynamicModulesLock);
+}
+static void unlockDynamicModules(void) {
+ LeaveCriticalSection(&DynamicModulesLock);
+}
+#else
+static pthread_once_t HipLoadedOnce = PTHREAD_ONCE_INIT;
+static pthread_mutex_t DynamicModulesLock = PTHREAD_MUTEX_INITIALIZER;
+static void lockDynamicModules(void) {
+ pthread_mutex_lock(&DynamicModulesLock);
+}
+static void unlockDynamicModules(void) {
+ pthread_mutex_unlock(&DynamicModulesLock);
+}
+#endif
+
+static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
+ const char *Target);
+
+static int isVerboseMode() {
+ static int IsVerbose = -1;
+ if (IsVerbose == -1)
+ IsVerbose = getenv("LLVM_PROFILE_VERBOSE") != nullptr;
+ return IsVerbose;
+}
+
+/* -------------------------------------------------------------------------- */
+/* Dynamic loading of HIP runtime symbols */
+/* -------------------------------------------------------------------------- */
+
+typedef int (*hipGetSymbolAddressTy)(void **, const void *);
+typedef int (*hipMemcpyTy)(void *, const void *, size_t, int);
+typedef int (*hipModuleGetGlobalTy)(void **, size_t *, void *, const char *);
+typedef int (*hipGetDeviceCountTy)(int *);
+typedef int (*hipGetDeviceTy)(int *);
+typedef int (*hipSetDeviceTy)(int);
+
+/* Minimal hipDeviceProp_t (HIP 6.x R0600): only gcnArchName at offset 1160
+ * is read. Padded to 4096 to tolerate ABI growth. */
+typedef struct {
+ char padding[1160];
+ char gcnArchName[256];
+ char tail_padding[2680];
+} HipDevicePropMinimal;
+typedef int (*hipGetDevicePropertiesTy)(HipDevicePropMinimal *, int);
+
+static hipGetSymbolAddressTy pHipGetSymbolAddress = nullptr;
+static hipMemcpyTy pHipMemcpy = nullptr;
+static hipModuleGetGlobalTy pHipModuleGetGlobal = nullptr;
+static hipGetDeviceCountTy pHipGetDeviceCount = nullptr;
+static hipGetDeviceTy pHipGetDevice = nullptr;
+static hipSetDeviceTy pHipSetDevice = nullptr;
+static hipGetDevicePropertiesTy pHipGetDeviceProperties = nullptr;
+
+static int NumDevices = 0;
+/* 256 matches hipDeviceProp_t::gcnArchName, the source field width. */
+static char (*DeviceArchNames)[256] = nullptr;
+
+/* -------------------------------------------------------------------------- */
+/* Device-to-host copies */
+/* Keep HIP-only to avoid an HSA dependency. */
+/* -------------------------------------------------------------------------- */
+
+static void doEnsureHipLoaded(void) {
+ if (!__interception::DynamicLoaderAvailable()) {
+ if (isVerboseMode())
+ PROF_NOTE("%s", "Dynamic library loading not available - "
+ "HIP profiling disabled\n");
+ return;
+ }
+
+#ifdef _WIN32
+ static const char HipLibName[] = "amdhip64.dll";
+#else
+ static const char HipLibName[] = "libamdhip64.so";
+#endif
+
+ void *Handle = __interception::OpenLibrary(HipLibName);
+ if (!Handle)
+ return;
+
+ pHipGetSymbolAddress = (hipGetSymbolAddressTy)__interception::LookupSymbol(
+ Handle, "hipGetSymbolAddress");
+ pHipMemcpy = (hipMemcpyTy)__interception::LookupSymbol(Handle, "hipMemcpy");
+ pHipModuleGetGlobal = (hipModuleGetGlobalTy)__interception::LookupSymbol(
+ Handle, "hipModuleGetGlobal");
+ pHipGetDeviceCount = (hipGetDeviceCountTy)__interception::LookupSymbol(
+ Handle, "hipGetDeviceCount");
+ pHipGetDevice =
+ (hipGetDeviceTy)__interception::LookupSymbol(Handle, "hipGetDevice");
+ pHipSetDevice =
+ (hipSetDeviceTy)__interception::LookupSymbol(Handle, "hipSetDevice");
+ pHipGetDeviceProperties =
+ (hipGetDevicePropertiesTy)__interception::LookupSymbol(
+ Handle, "hipGetDevicePropertiesR0600");
+ if (!pHipGetDeviceProperties)
+ pHipGetDeviceProperties =
+ (hipGetDevicePropertiesTy)__interception::LookupSymbol(
+ Handle, "hipGetDeviceProperties");
+
+ if (pHipGetDeviceCount && pHipGetDeviceProperties) {
+ int Count = 0;
+ if (pHipGetDeviceCount(&Count) == 0 && Count > 0) {
+ DeviceArchNames = (char (*)[256])calloc(Count, sizeof(*DeviceArchNames));
+ if (!DeviceArchNames) {
+ PROF_ERR("%s\n", "failed to allocate device arch name table");
+ return;
+ }
+ HipDevicePropMinimal Prop;
+ for (int i = 0; i < Count; ++i) {
+ __builtin_memset(&Prop, 0, sizeof(Prop));
+ if (pHipGetDeviceProperties(&Prop, i) == 0) {
+ strncpy(DeviceArchNames[i], Prop.gcnArchName,
+ sizeof(DeviceArchNames[i]) - 1);
+ DeviceArchNames[i][sizeof(DeviceArchNames[i]) - 1] = '\0';
+ if (isVerboseMode())
+ PROF_NOTE("Device %d arch: %s\n", i, DeviceArchNames[i]);
+ }
+ }
+ NumDevices = Count;
+ }
+ }
+}
+
+#ifdef _WIN32
+static BOOL CALLBACK ensureHipLoadedCb(PINIT_ONCE, PVOID, PVOID *) {
+ doEnsureHipLoaded();
+ return TRUE;
+}
+#endif
+
+static void ensureHipLoaded(void) {
+#ifdef _WIN32
+ InitOnceExecuteOnce(&HipLoadedOnce, ensureHipLoadedCb, NULL, NULL);
+#else
+ pthread_once(&HipLoadedOnce, doEnsureHipLoaded);
+#endif
+}
+
+/* -------------------------------------------------------------------------- */
+/* Public wrappers that forward to the loaded HIP symbols */
+/* -------------------------------------------------------------------------- */
+
+static int hipGetSymbolAddress(void **devPtr, const void *symbol) {
+ ensureHipLoaded();
+ return pHipGetSymbolAddress ? pHipGetSymbolAddress(devPtr, symbol) : -1;
+}
+
+static int hipMemcpy(void *dest, const void *src, size_t len,
+ int kind /*2=DToH*/) {
+ ensureHipLoaded();
+ return pHipMemcpy ? pHipMemcpy(dest, src, len, kind) : -1;
+}
+
+/* Device section symbols must be registered with CLR first; otherwise
+ * hipMemcpy may take a CPU path and crash. */
+static int memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) {
+ return hipMemcpy(Dst, Src, Size, 2 /* DToH */);
+}
+
+static int hipModuleGetGlobal(void **DevPtr, size_t *Bytes, void *Module,
+ const char *Name) {
+ ensureHipLoaded();
+ return pHipModuleGetGlobal ? pHipModuleGetGlobal(DevPtr, Bytes, Module, Name)
+ : -1;
+}
+
+static int hipGetDevice(int *DeviceId) {
+ ensureHipLoaded();
+ return pHipGetDevice ? pHipGetDevice(DeviceId) : -1;
+}
+
+static int hipSetDevice(int DeviceId) {
+ ensureHipLoaded();
+ return pHipSetDevice ? pHipSetDevice(DeviceId) : -1;
+}
+
+static const char *getDeviceArchName(int DeviceId) {
+ if (DeviceId < 0 || DeviceId >= NumDevices || !DeviceArchNames[DeviceId][0])
+ return "amdgpu";
+ return DeviceArchNames[DeviceId];
+}
+
+/* -------------------------------------------------------------------------- */
+/* Dynamic module tracking */
+/* -------------------------------------------------------------------------- */
+
+/* Per-TU profile entry inside a dynamic module.
+ * A single dynamic module may contain multiple TUs (e.g. -fgpu-rdc). */
+typedef struct {
+ void *DeviceVar; /* device address of __llvm_profile_sections_<CUID> */
+ int Processed; /* 0 = not yet collected, 1 = data already copied */
+} OffloadDynamicTUInfo;
+
+/* One entry per hipModuleLoad call. */
+typedef struct {
+ void *ModulePtr; /* hipModule_t handle */
+ OffloadDynamicTUInfo *TUs; /* array of per-TU entries */
+ int NumTUs;
+ int CapTUs;
+} OffloadDynamicModuleInfo;
+
+static OffloadDynamicModuleInfo *DynamicModules = nullptr;
+static int NumDynamicModules = 0;
+static int CapDynamicModules = 0;
+
+/* -------------------------------------------------------------------------- */
+/* ELF symbol enumeration (manual parse: compiler-rt cannot link LLVM Support)
+ */
+/* -------------------------------------------------------------------------- */
+
+#if __has_include(<elf.h>)
+#include <elf.h>
+
+/* Callback invoked for every matching symbol name found in the ELF image.
+ * Return 0 to continue iteration, non-zero to stop. */
+typedef int (*SymbolCallback)(const char *Name, void *UserData);
+
+/* If Image is a clang offload bundle, return a pointer to the first embedded
+ * ELF. Returns Image if not a bundle, nullptr if a bundle holds no ELF. */
+static const void *unwrapOffloadBundle(const void *Image) {
+ static const char BundleMagic[] = "__CLANG_OFFLOAD_BUNDLE__";
+ if (memcmp(Image, BundleMagic, sizeof(BundleMagic) - 1) != 0)
+ return Image; /* Not a bundle, return as-is. */
+
+ const char *Buf = (const char *)Image;
+ uint64_t NumEntries;
+ __builtin_memcpy(&NumEntries, Buf + sizeof(BundleMagic) - 1,
+ sizeof(uint64_t));
+
+ /* Walk the entry table (starts at offset 32). */
+ const char *Cursor = Buf + 32;
+ for (uint64_t I = 0; I < NumEntries; ++I) {
+ uint64_t EntryOffset, EntrySize, IDSize;
+ __builtin_memcpy(&EntryOffset, Cursor, sizeof(EntryOffset));
+ Cursor += sizeof(EntryOffset);
+ __builtin_memcpy(&EntrySize, Cursor, sizeof(EntrySize));
+ Cursor += sizeof(EntrySize);
+ __builtin_memcpy(&IDSize, Cursor, sizeof(IDSize));
+ Cursor += sizeof(IDSize);
+ Cursor += IDSize; /* skip entry ID */
+
+ if (EntrySize >= sizeof(Elf64_Ehdr)) {
+ const Elf64_Ehdr *E = (const Elf64_Ehdr *)(Buf + EntryOffset);
+ if (E->e_ident[EI_MAG0] == ELFMAG0 && E->e_ident[EI_MAG1] == ELFMAG1 &&
+ E->e_ident[EI_MAG2] == ELFMAG2 && E->e_ident[EI_MAG3] == ELFMAG3) {
+ return (const void *)(Buf + EntryOffset);
+ }
+ }
+ }
+
+ PROF_WARN("%s", "offload bundle contains no valid ELF entries\n");
+ return nullptr;
+}
+
+/* Invoke CB for every global symbol in Image (an AMDGPU ELF or offload bundle)
+ * whose name starts with PREFIX. Image may be null. */
+static void enumerateElfSymbols(const void *Image, const char *Prefix,
+ SymbolCallback CB, void *UserData) {
+ if (!Image)
+ return;
+
+ Image = unwrapOffloadBundle(Image);
+ if (!Image)
+ return;
+
+ const Elf64_Ehdr *Ehdr = (const Elf64_Ehdr *)Image;
+ if (Ehdr->e_ident[EI_MAG0] != ELFMAG0 || Ehdr->e_ident[EI_MAG1] != ELFMAG1 ||
+ Ehdr->e_ident[EI_MAG2] != ELFMAG2 || Ehdr->e_ident[EI_MAG3] != ELFMAG3) {
+ if (isVerboseMode())
+ PROF_NOTE("%s", "Image is not a valid ELF, skipping enumeration\n");
+ return;
+ }
+
+ size_t PrefixLen = strlen(Prefix);
+ const char *Base = (const char *)Image;
+ const Elf64_Shdr *Shdrs = (const Elf64_Shdr *)(Base + Ehdr->e_shoff);
+
+ for (int i = 0; i < Ehdr->e_shnum; ++i) {
+ if (Shdrs[i].sh_type != SHT_SYMTAB)
+ continue;
+
+ const Elf64_Sym *Syms = (const Elf64_Sym *)(Base + Shdrs[i].sh_offset);
+ int NumSyms = Shdrs[i].sh_size / sizeof(Elf64_Sym);
+ /* String table is the section referenced by sh_link. */
+ const char *StrTab = Base + Shdrs[Shdrs[i].sh_link].sh_offset;
+
+ for (int j = 0; j < NumSyms; ++j) {
+ if (Syms[j].st_name == 0)
+ continue;
+ const char *Name = StrTab + Syms[j].st_name;
+ if (strncmp(Name, Prefix, PrefixLen) == 0) {
+ if (CB(Name, UserData))
+ return;
+ }
+ }
+ }
+}
+
+/* State passed through the enumeration callback. */
+typedef struct {
+ void *Module; /* hipModule_t */
+ OffloadDynamicModuleInfo *ModInfo;
+} EnumState;
+
+/* Register one __llvm_profile_sections_<CUID> symbol on the module entry.
+ * hipModuleGetGlobal also registers the device address with CLR so hipMemcpy
+ * can copy from it later. */
+static int registerPrfSymbol(const char *Name, void *UserData) {
+ EnumState *S = (EnumState *)UserData;
+ OffloadDynamicModuleInfo *MI = S->ModInfo;
+
+ /* The symbol is the per-TU sections struct itself, not a pointer
+ * indirection, so this address is the hipMemcpy source. */
+ void *DeviceVar = nullptr;
+ size_t Bytes = 0;
+ if (hipModuleGetGlobal(&DeviceVar, &Bytes, S->Module, Name) != 0) {
+ PROF_WARN("failed to get symbol %s for module %p\n", Name, S->Module);
+ return 0; /* continue */
+ }
+
+ if (MI->NumTUs >= MI->CapTUs) {
+ int NewCap = MI->CapTUs ? MI->CapTUs * 2 : 4;
+ OffloadDynamicTUInfo *New = (OffloadDynamicTUInfo *)realloc(
+ MI->TUs, NewCap * sizeof(OffloadDynamicTUInfo));
+ if (!New) {
+ PROF_ERR("%s\n", "failed to grow TU array");
+ return 0;
+ }
+ MI->TUs = New;
+ MI->CapTUs = NewCap;
+ }
+ OffloadDynamicTUInfo *TU = &MI->TUs[MI->NumTUs++];
+ TU->DeviceVar = DeviceVar;
+ TU->Processed = 0;
+
+ (void)Name;
+ return 0; /* continue enumeration */
+}
+
+#endif /* __has_include(<elf.h>) */
+
+/* -------------------------------------------------------------------------- */
+/* Registration / un-registration helpers */
+/* -------------------------------------------------------------------------- */
+
+extern "C" void
+__llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
+ const void *Image) {
+ if (ModuleLoadRc)
+ return;
+
+ lockDynamicModules();
+
+ if (isVerboseMode())
+ PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n",
+ NumDynamicModules, ModuleLoadRc, *Ptr, Image);
+
+ if (NumDynamicModules >= CapDynamicModules) {
+ int NewCap = CapDynamicModules ? CapDynamicModules * 2 : 64;
+ OffloadDynamicModuleInfo *New = (OffloadDynam...
[truncated]
|
…ined library Move the host-side ROCm/HIP device profile collection runtime out of clang_rt.profile into a separate, opt-in library clang_rt.profile_rocm, and make that library self-contained. The ROCm collection runtime (InstrProfilingPlatformROCm.cpp) depends on the sanitizer interception + sanitizer_common object libraries, which are built with the dynamic CRT (/MD) on Windows. Merging them into clang_rt.profile forced the whole profile archive to /MD and broke existing static-CRT (/MT) consumers, which is why llvm#177665 was reverted. clang_rt.profile_rocm is a superset of clang_rt.profile: all the base profile sources plus InstrProfilingPlatformROCm.cpp and the interceptor object libs, with -DCOMPILER_RT_BUILD_PROFILE_ROCM=1. Because it contains the whole base runtime it is built with a single CRT model throughout (/MD on Windows), so there is no CRT mixing inside the archive. clang_rt.profile is left exactly as before: /MT on Windows, no ROCm dependency, no interceptor merge. The driver links clang_rt.profile_rocm ahead of clang_rt.profile for HIP device PGO. Being a superset linked first, it resolves all profile symbols, so the base archive contributes no objects and the two CRT models are never mixed in the host image (verified: linking both produces no duplicate symbols and the base archive stays inert). clang_rt.profile_rocm is gated on COMPILER_RT_BUILD_PROFILE_ROCM, which now defaults off on all platforms; a default build is never affected.
yxsamliu
force-pushed
the
amd/dev/yaxunl/offload-pgo-runtime-main
branch
from
June 4, 2026 15:45
ff8aa52 to
7948ad1
Compare
Contributor
Author
|
Updated: clang_rt.profile_rocm is now a self-contained superset of clang_rt.profile (all base sources plus the ROCm collector and interceptor deps), built entirely /MD on Windows. This keeps a single CRT model within the archive and leaves clang_rt.profile completely untouched (/MT, no ROCm). The driver links profile_rocm ahead of the base profile, which then stays inert, so the two CRT models are never mixed in the host image. Verified by linking a profiling program against both archives: no duplicate symbols, base archive inert. |
jhuber6
approved these changes
Jun 4, 2026
chapuni
added a commit
that referenced
this pull request
Jun 5, 2026
This has been intrdocued in #201606.
This was referenced Jun 6, 2026
lfmeadow
added a commit
to ROCm/llvm-project
that referenced
this pull request
Jun 9, 2026
…eland amd-staging now carries the relanded HIP offload PGO runtime as the opt-in clang_rt.profile_rocm library (llvm#201606), but lacks the host-link wiring (llvm#201607, still open upstream), so a HIP+PGO host link leaves the __llvm_profile_offload_* symbols unresolved. This restarts the device-drain work cleanly on top of that reland instead of rebasing the prior branch. Drain runtime (split by platform, mutually exclusive guards so exactly one provides the __llvm_profile_offload_* / __llvm_profile_hip_collect_device_data symbols): - InstrProfilingPlatformROCm.cpp: Linux HSA introspection. Walks every GPU agent's loaded code objects at exit, so it drains device counters even for device modules with no host-side shadow (RCCL-style synthesized kernels) and for kernels on a non-current device. - InstrProfilingPlatformROCmWindows.cpp: the reland host-shadow drain, unchanged, guarded to _WIN32 (Windows has no HSA introspection). - CMakeLists.txt wires both into clang_rt.profile_rocm. Driver host-link wiring (net-new part of llvm#201607, adapted to amd-staging): - Linux.cpp / MSVC.cpp link clang_rt.profile_rocm ahead of the base profile for HIP host links when profiling is requested. - hip-profile-rocm-runtime.hip covers the link ordering. Standalone (non-TheRock) build + test recipe: - compiler-rt/test/profile/device-pgo/{toolchain-cache.cmake,build.sh,README.md} build the toolchain + clang_rt.profile_rocm + the amdgcn device profile runtime in one cmake configure (LLVM_RUNTIME_TARGETS=default;amdgcn-amd-amdhsa). - run_gpu_tests.py (lit-lite runner) derives the multi-device feature from the runtime-visible GPU count via amdgpu-arch; GPU/ and AMDGPU/ .hip tests use Linux HSA CHECK strings and mark agent-walk-only cases UNSUPPORTED: windows. Co-authored-by: Cursor <cursoragent@cursor.com>
carlobertolli
pushed a commit
to carlobertolli/llvm-project
that referenced
this pull request
Jun 11, 2026
…lvm#201606) This mostly relands the compiler-rt part of llvm#177665 (approved and merged, then reverted in llvm#201416). The first commit restores it as merged. It was reverted because of a Windows problem: the ROCm runtime needs the sanitizer interception library, which is built /MD on Windows, so putting it in clang_rt.profile forced that library to /MD and broke users linking it with the static CRT (/MT). The second commit fixes this by building the ROCm support as a separate, opt-in library clang_rt.profile_rocm, a /MD superset of clang_rt.profile. The base library is left unchanged (/MT, no ROCm). The driver links clang_rt.profile_rocm first, so it resolves all profile symbols and the base library stays inert. clang_rt.profile_rocm is off by default. The compiler-side change and driver wiring are in a separate PR.
carlobertolli
pushed a commit
to carlobertolli/llvm-project
that referenced
this pull request
Jun 11, 2026
This has been intrdocued in llvm#201606.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This mostly relands the compiler-rt part of #177665 (approved and merged, then reverted in #201416). The first commit restores it as merged.
It was reverted because of a Windows problem: the ROCm runtime needs the sanitizer interception library, which is built /MD on Windows, so putting it in clang_rt.profile forced that library to /MD and broke users linking it with the static CRT (/MT).
The second commit fixes this by building the ROCm support as a separate, opt-in library clang_rt.profile_rocm, a /MD superset of clang_rt.profile. The base library is left unchanged (/MT, no ROCm). The driver links clang_rt.profile_rocm first, so it resolves all profile symbols and the base library stays inert.
clang_rt.profile_rocm is off by default. The compiler-side change and driver wiring are in a separate PR.