[clang] [compiler-rt] [llvm] [PGO][AMDGPU] Add basic HIP offload PGO support (PR #177665)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed May 20 07:04:54 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/177665
>From 4897c125760a5eb59ef10c6e08a0a2e7d3272a85 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 10 Apr 2026 15:01:37 -0400
Subject: [PATCH 01/13] [PGO][AMDGPU] Add basic HIP offload PGO support
Add the ROCm-side runtime and compiler support needed to collect AMDGPU
device profiles from HIP programs, merge them with host profiles, and use
them for PGO. This includes InstrProfilingPlatformROCm for device data
collection and host write paths, AMDGPU lowering to
__llvm_profile_instrument_gpu with regular counters, and disabling GPU
indirect-call value profiling where unsupported.
Also adjust compiler-rt profile build flags (-fno-exceptions, MSVC
MultiThreadedDLL when merging interception/sanitizer objects into the profile
archive) and align Windows lit configuration for Profile-* tests and
ASan+llvm-coverage with the profile runtime's /MD CRT expectations.
---
compiler-rt/lib/profile/CMakeLists.txt | 31 +-
compiler-rt/lib/profile/InstrProfilingFile.c | 6 +
.../profile/InstrProfilingPlatformROCm.cpp | 830 ++++++++++++++++++
.../Instrumentation/InstrProfiling.cpp | 462 +++++++++-
.../Instrumentation/PGOInstrumentation.cpp | 11 +-
.../amdgpu-contiguous-counters.ll | 34 +
.../InstrProfiling/amdgpu-instrumentation.ll | 32 +
.../InstrProfiling/gpu-weak.ll | 36 +
.../amdgpu-disable-value-profiling.ll | 22 +
9 files changed, 1426 insertions(+), 38 deletions(-)
create mode 100644 compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
create mode 100644 llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
create mode 100644 llvm/test/Instrumentation/InstrProfiling/gpu-weak.ll
create mode 100644 llvm/test/Transforms/PGOProfile/amdgpu-disable-value-profiling.ll
diff --git a/compiler-rt/lib/profile/CMakeLists.txt b/compiler-rt/lib/profile/CMakeLists.txt
index 8d9a773412a22..36e6b5a6a21c1 100644
--- a/compiler-rt/lib/profile/CMakeLists.txt
+++ b/compiler-rt/lib/profile/CMakeLists.txt
@@ -89,6 +89,7 @@ if (NOT COMPILER_RT_PROFILE_BAREMETAL)
list(APPEND PROFILE_SOURCES
GCDAProfiling.c
InstrProfilingFile.c
+ InstrProfilingPlatformROCm.cpp
InstrProfilingRuntime.cpp
InstrProfilingUtil.c
InstrProfilingValue.c
@@ -155,6 +156,17 @@ if(COMPILER_RT_PROFILE_BAREMETAL)
-DCOMPILER_RT_PROFILE_BAREMETAL=1)
endif()
+set(PROFILE_OBJECT_LIBS)
+if(COMPILER_RT_HAS_INTERCEPTION AND NOT COMPILER_RT_PROFILE_BAREMETAL)
+ # RTInterception references __sanitizer_internal_{memcpy,memset,memmove} and other
+ # sanitizer_common symbols; merge the same object libs as clang_rt.cfi (without
+ # coverage/symbolizer) so -fprofile-instr-generate links stay self-contained.
+ list(APPEND PROFILE_OBJECT_LIBS
+ RTInterception
+ RTSanitizerCommon
+ RTSanitizerCommonLibc)
+endif()
+
if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx")
append_list_if(COMPILER_RT_HAS_FFREESTANDING_FLAG -ffreestanding EXTRA_FLAGS)
append_list_if(COMPILER_RT_HAS_NOGPULIB_FLAG -nogpulib EXTRA_FLAGS)
@@ -168,13 +180,24 @@ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx")
endif()
if(MSVC)
- # profile historically has only been supported with the static runtime
- # on windows
- set(CMAKE_MSVC_RUNTIME_LIBRARY MultiThreaded)
+ # profile historically used the static CRT (/MT). When we merge RTInterception and
+ # RTSanitizerCommon (same object libs as clang_rt.cfi on ELF), those targets are
+ # built with MultiThreadedDLL (/MD) — see interception/CMakeLists.txt and
+ # sanitizer_common/CMakeLists.txt. Mixing /MD objects into a /MT libclang_rt.profile
+ # yields LNK2019 (__imp__stricmp from interception_win.cpp) and LNK4098 in Profile-*.
+ if(COMPILER_RT_HAS_INTERCEPTION AND NOT COMPILER_RT_PROFILE_BAREMETAL)
+ set(CMAKE_MSVC_RUNTIME_LIBRARY MultiThreadedDLL)
+ else()
+ set(CMAKE_MSVC_RUNTIME_LIBRARY MultiThreaded)
+ endif()
endif()
# We don't use the C++ Standard Library here, so avoid including it by mistake.
append_list_if(COMPILER_RT_HAS_NOSTDINCXX_FLAG -nostdinc++ EXTRA_FLAGS)
+# C++ profile sources (e.g. InstrProfilingPlatformROCm.cpp) must not emit exception
+# personality symbols: host libclang_rt.profile.a is linked from C code and from C++
+# tests that do not pull in __gxx_personality_v0 (Profile-* / premerge).
+append_list_if(COMPILER_RT_HAS_FNO_EXCEPTIONS_FLAG -fno-exceptions EXTRA_FLAGS)
# XRay uses C++ standard library headers.
string(REGEX REPLACE "-?-stdlib=[a-zA-Z+]*" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
@@ -200,6 +223,7 @@ if(APPLE)
STATIC
OS ${PROFILE_SUPPORTED_OS}
ARCHS ${PROFILE_SUPPORTED_ARCH}
+ OBJECT_LIBS ${PROFILE_OBJECT_LIBS}
CFLAGS ${EXTRA_FLAGS}
SOURCES ${PROFILE_SOURCES}
ADDITIONAL_HEADERS ${PROFILE_HEADERS}
@@ -209,6 +233,7 @@ else()
add_compiler_rt_runtime(clang_rt.profile
STATIC
ARCHS ${PROFILE_SUPPORTED_ARCH}
+ OBJECT_LIBS ${PROFILE_OBJECT_LIBS}
CFLAGS ${EXTRA_FLAGS}
SOURCES ${PROFILE_SOURCES}
ADDITIONAL_HEADERS ${PROFILE_HEADERS}
diff --git a/compiler-rt/lib/profile/InstrProfilingFile.c b/compiler-rt/lib/profile/InstrProfilingFile.c
index 71127b05aafb8..dae55a5fc0016 100644
--- a/compiler-rt/lib/profile/InstrProfilingFile.c
+++ b/compiler-rt/lib/profile/InstrProfilingFile.c
@@ -41,6 +41,10 @@
#include "InstrProfilingPort.h"
#include "InstrProfilingUtil.h"
+/* HIP / offload collection hook implemented in InstrProfilingPlatformROCm.c.
+ * It is a no-op when no offload profile data was registered. */
+extern int __llvm_profile_hip_collect_device_data(void);
+
/* From where is profile name specified.
* The order the enumerators define their
* precedence. Re-order them may lead to
@@ -1198,6 +1202,8 @@ int __llvm_profile_write_file(void) {
if (rc)
PROF_ERR("Failed to write file \"%s\": %s\n", Filename, strerror(errno));
+ __llvm_profile_hip_collect_device_data();
+
// Restore SIGKILL.
if (PDeathSig == 1)
lprofRestoreSigKill();
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
new file mode 100644
index 0000000000000..c0b1732c8d75b
--- /dev/null
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -0,0 +1,830 @@
+//===- 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"
+}
+
+// interception.h pulls in sanitizer_internal_defs.h, which normally includes
+// sanitizer_redefine_builtins.h. That uses inline asm to alias
+// memcpy/memmove/memset to __sanitizer_internal_* (see sanitizer_libc.cpp in
+// sanitizer_common). The instrumented *host* link for HIP (-fprofile-generate)
+// only pulls in libclang_rt.profile.a, not the full sanitizer_common objects
+// that define those symbols, so we get undefined references at link time. This
+// TU does not need the sanitizer builtin redirect; keep using libc
+// memcpy/memset.
+#define SANITIZER_COMMON_NO_REDEFINE_BUILTINS 1
+#include "interception/interception.h"
+#undef SANITIZER_COMMON_NO_REDEFINE_BUILTINS
+// 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>
+
+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);
+
+/* hipDeviceProp_t layout for HIP 6.x+ (R0600).
+ * We only need gcnArchName at offset 1160. Pad to 4096 to safely
+ * accommodate future struct growth without recompilation. */
+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;
+
+#define MAX_DEVICES 16
+static int NumDevices = 0;
+static char DeviceArchNames[MAX_DEVICES][256];
+
+/* -------------------------------------------------------------------------- */
+/* Device-to-host copies */
+/* Keep HIP-only to avoid an HSA dependency. */
+/* -------------------------------------------------------------------------- */
+
+static void EnsureHipLoaded(void) {
+ static int Initialized = 0;
+ if (Initialized)
+ return;
+ Initialized = 1;
+
+ 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) {
+ if (Count > MAX_DEVICES)
+ Count = MAX_DEVICES;
+ HipDevicePropMinimal Prop;
+ for (int i = 0; i < Count; ++i) {
+ 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;
+ }
+ }
+}
+
+/* -------------------------------------------------------------------------- */
+/* 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;
+}
+
+/* Copy from device to host using HIP.
+ * This requires that the device section symbols are registered with CLR,
+ * otherwise hipMemcpy may attempt 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_offload_prf_<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 */
+/* */
+/* AMDGPU code objects are always ELF. We use manual parsing because this */
+/* is compiler-rt (standalone C runtime) and cannot link against LLVM's C++ */
+/* Support libraries. */
+/* -------------------------------------------------------------------------- */
+
+#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 (__CLANG_OFFLOAD_BUNDLE__), find the
+ * first embedded code object that is a valid ELF and return a pointer to it.
+ * Otherwise return Image unchanged. Returns nullptr if no ELF is found. */
+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;
+ 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;
+ memcpy(&EntryOffset, Cursor, 8);
+ Cursor += 8;
+ memcpy(&EntrySize, Cursor, 8);
+ Cursor += 8;
+ memcpy(&IDSize, Cursor, 8);
+ Cursor += 8;
+ /* Skip the entry ID string. */
+ Cursor += IDSize;
+
+ /* Check if this entry contains an ELF. */
+ 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;
+}
+
+/* Parse an AMDGPU code-object ELF and invoke CB for every global symbol whose
+ * name starts with PREFIX. Image may be nullptr (e.g. hipModuleLoad from file)
+ * or a clang offload bundle containing an ELF;
+ * in that case the function unwraps the bundle first. */
+static void EnumerateElfSymbols(const void *Image, const char *Prefix,
+ SymbolCallback CB, void *UserData) {
+ if (!Image)
+ return;
+
+ /* Handle clang offload bundle wrapping. */
+ 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;
+
+/* Grow the TU array inside a module entry and register one __llvm_offload_prf_*
+ * symbol. Also pre-registers the corresponding per-TU section symbols with CLR
+ * (needed so hipMemcpy can copy from those device addresses later). */
+static int RegisterPrfSymbol(const char *Name, void *UserData) {
+ EnumState *S = (EnumState *)UserData;
+ OffloadDynamicModuleInfo *MI = S->ModInfo;
+
+ /* Look up the per-TU pointer variable, then dereference to get the
+ * address of __llvm_profile_sections. */
+ void *DevicePtrVar = nullptr;
+ size_t Bytes = 0;
+ if (hipModuleGetGlobal(&DevicePtrVar, &Bytes, S->Module, Name) != 0) {
+ PROF_WARN("failed to get symbol %s for module %p\n", Name, S->Module);
+ return 0; /* continue */
+ }
+ void *DeviceVar = nullptr;
+ if (hipMemcpy(&DeviceVar, DevicePtrVar, sizeof(void *), 2 /*DToH*/) != 0) {
+ PROF_WARN("failed to read sections pointer for %s\n", Name);
+ return 0;
+ }
+
+ /* Grow TU array if needed. */
+ 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; /* CUID suffix available for future per-TU section lookup */
+
+ 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 (IsVerboseMode())
+ PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n",
+ NumDynamicModules, ModuleLoadRc, *Ptr, Image);
+
+ if (ModuleLoadRc)
+ return;
+
+ if (NumDynamicModules >= CapDynamicModules) {
+ int NewCap = CapDynamicModules ? CapDynamicModules * 2 : 64;
+ OffloadDynamicModuleInfo *New = (OffloadDynamicModuleInfo *)realloc(
+ DynamicModules, NewCap * sizeof(OffloadDynamicModuleInfo));
+ if (!New)
+ return;
+ DynamicModules = New;
+ CapDynamicModules = NewCap;
+ }
+
+ OffloadDynamicModuleInfo *MI = &DynamicModules[NumDynamicModules++];
+ MI->ModulePtr = *Ptr;
+ MI->TUs = nullptr;
+ MI->NumTUs = 0;
+ MI->CapTUs = 0;
+
+ /* Enumerate all __llvm_offload_prf_<CUID> symbols in the ELF image.
+ * For each one, look it up via hipModuleGetGlobal (which also registers
+ * the device address with CLR for later hipMemcpy) and store the entry.
+ *
+ * ELF parsing requires <elf.h>. On platforms without it, dynamic module
+ * profiling is not yet supported. */
+#if __has_include(<elf.h>)
+ EnumState State = {*Ptr, MI};
+ EnumerateElfSymbols(Image, "__llvm_offload_prf_", RegisterPrfSymbol, &State);
+#else
+ (void)Image;
+ if (IsVerboseMode())
+ PROF_NOTE("%s",
+ "Dynamic module profiling not supported on this platform\n");
+#endif
+
+ if (MI->NumTUs == 0) {
+ PROF_WARN("no __llvm_offload_prf_* symbols found in module %p\n", *Ptr);
+ } else if (IsVerboseMode()) {
+ PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs);
+ }
+}
+
+extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
+ for (int i = 0; i < NumDynamicModules; ++i) {
+ OffloadDynamicModuleInfo *MI = &DynamicModules[i];
+
+ if (MI->ModulePtr != Ptr)
+ continue;
+
+ if (IsVerboseMode())
+ PROF_NOTE("Unregistering module %p (%d TUs)\n", MI->ModulePtr,
+ MI->NumTUs);
+
+ /* Process every TU in this module. */
+ for (int t = 0; t < MI->NumTUs; ++t) {
+ OffloadDynamicTUInfo *TU = &MI->TUs[t];
+ if (TU->Processed) {
+ if (IsVerboseMode())
+ PROF_NOTE("Module %p TU %d already processed, skipping\n", Ptr, t);
+ continue;
+ }
+ /* Use a globally unique index as TU index for the output filename. */
+ int TUIndex = i * 1000 + t;
+ if (TU->DeviceVar) {
+ int CurDev = 0;
+ hipGetDevice(&CurDev);
+ const char *ArchName = getDeviceArchName(CurDev);
+ if (ProcessDeviceOffloadPrf(TU->DeviceVar, TUIndex, ArchName) == 0)
+ TU->Processed = 1;
+ else
+ PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr,
+ t);
+ }
+ }
+ return;
+ }
+
+ if (IsVerboseMode())
+ PROF_WARN("unregister called for unknown module %p\n", Ptr);
+}
+
+/* Grow a void* array, doubling capacity (or starting at InitCap). */
+static int GrowPtrArray(void ***Arr, int *Num, int *Cap, int InitCap) {
+ if (*Num < *Cap)
+ return 0;
+ int NewCap = *Cap ? *Cap * 2 : InitCap;
+ void **New = (void **)realloc(*Arr, NewCap * sizeof(void *));
+ if (!New)
+ return -1;
+ *Arr = New;
+ *Cap = NewCap;
+ return 0;
+}
+
+static void **OffloadShadowVariables = nullptr;
+static int NumShadowVariables = 0;
+static int CapShadowVariables = 0;
+
+extern "C" void __llvm_profile_offload_register_shadow_variable(void *ptr) {
+ if (GrowPtrArray(&OffloadShadowVariables, &NumShadowVariables,
+ &CapShadowVariables, 64))
+ return;
+ OffloadShadowVariables[NumShadowVariables++] = ptr;
+}
+
+static void **OffloadSectionShadowVariables = nullptr;
+static int NumSectionShadowVariables = 0;
+static int CapSectionShadowVariables = 0;
+
+extern "C" void
+__llvm_profile_offload_register_section_shadow_variable(void *ptr) {
+ if (GrowPtrArray(&OffloadSectionShadowVariables, &NumSectionShadowVariables,
+ &CapSectionShadowVariables, 64))
+ return;
+ OffloadSectionShadowVariables[NumSectionShadowVariables++] = ptr;
+}
+
+// Free host-side copies of device sections on error or success. Factored out
+// so we can return early: C++ forbids goto past initializations of automatic
+// locals declared later in this function (e.g. const uint64_t NumData).
+// Callers pass nullptr for reused (cached) sections so we only free malloc'd
+// buffers; free(nullptr) is a no-op (C/C++).
+static void freeCopiedHostSections(char *HostCountersBegin, char *HostDataBegin,
+ char *HostNamesBegin) {
+ free(HostCountersBegin);
+ free(HostDataBegin);
+ free(HostNamesBegin);
+}
+
+namespace {
+
+struct CopiedHostSectionsCleanup {
+ char *Counters;
+ char *Data;
+ char *Names;
+ int CntsReused;
+ int DataReused;
+ int NamesReused;
+
+ CopiedHostSectionsCleanup(char *C, char *D, char *N, int CR, int DR, int NR)
+ : Counters(C), Data(D), Names(N), CntsReused(CR), DataReused(DR),
+ NamesReused(NR) {}
+
+ ~CopiedHostSectionsCleanup() {
+ freeCopiedHostSections(CntsReused ? nullptr : Counters,
+ DataReused ? nullptr : Data,
+ NamesReused ? nullptr : Names);
+ }
+
+ CopiedHostSectionsCleanup(const CopiedHostSectionsCleanup &) = delete;
+ CopiedHostSectionsCleanup &
+ operator=(const CopiedHostSectionsCleanup &) = delete;
+};
+
+struct MallocBufferCleanup {
+ void *Ptr;
+ explicit MallocBufferCleanup(void *P) : Ptr(P) {}
+ ~MallocBufferCleanup() { free(Ptr); }
+ MallocBufferCleanup(const MallocBufferCleanup &) = delete;
+ MallocBufferCleanup &operator=(const MallocBufferCleanup &) = delete;
+ char *get() const { return static_cast<char *>(Ptr); }
+};
+
+} // namespace
+
+static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
+ const char *Target) {
+ __llvm_profile_gpu_sections HostSections;
+
+ if (hipMemcpy(&HostSections, DeviceOffloadPrf, sizeof(HostSections),
+ 2 /*DToH*/) != 0) {
+ PROF_ERR("%s\n", "failed to copy offload prf structure from device");
+ return -1;
+ }
+
+ const void *DevCntsBegin = HostSections.CountersStart;
+ const void *DevDataBegin = HostSections.DataStart;
+ const void *DevNamesBegin = HostSections.NamesStart;
+ const void *DevCntsEnd = HostSections.CountersStop;
+ const void *DevDataEnd = HostSections.DataStop;
+ const void *DevNamesEnd = HostSections.NamesStop;
+
+ size_t CountersSize = (const char *)DevCntsEnd - (const char *)DevCntsBegin;
+ size_t DataSize = (const char *)DevDataEnd - (const char *)DevDataBegin;
+ size_t NamesSize = (const char *)DevNamesEnd - (const char *)DevNamesBegin;
+
+ if (IsVerboseMode())
+ PROF_NOTE("Section pointers: Cnts=[%p,%p]=%zu Data=[%p,%p]=%zu "
+ "Names=[%p,%p]=%zu\n",
+ DevCntsBegin, DevCntsEnd, CountersSize, DevDataBegin, DevDataEnd,
+ DataSize, DevNamesBegin, DevNamesEnd, NamesSize);
+
+ if (CountersSize == 0 || DataSize == 0)
+ return 0;
+
+ int ret = -1;
+ int NamesReused = 0, CntsReused = 0, DataReused = 0;
+
+ char *HostDataBegin = nullptr;
+ char *HostCountersBegin = nullptr;
+ char *HostNamesBegin = nullptr;
+
+ /* Sections using linker-defined __start_/__stop_ bounds are shared across
+ TU structs in RDC mode. Deduplicate by caching the last copied range. */
+ static const void *CachedDevNamesBegin = nullptr;
+ static char *CachedHostNames = nullptr;
+ static size_t CachedNamesSize = 0;
+
+ static const void *CachedDevCntsBegin = nullptr;
+ static char *CachedHostCnts = nullptr;
+ static size_t CachedCntsSize = 0;
+
+ static const void *CachedDevDataBegin = nullptr;
+ static char *CachedHostData = nullptr;
+ static size_t CachedDataSize = 0;
+
+ if (CountersSize > 0 && DevCntsBegin == CachedDevCntsBegin &&
+ CountersSize == CachedCntsSize) {
+ HostCountersBegin = CachedHostCnts;
+ CntsReused = 1;
+ if (IsVerboseMode())
+ PROF_NOTE("Reusing cached counters section (%zu bytes)\n", CountersSize);
+ } else if (CountersSize > 0) {
+ HostCountersBegin = (char *)malloc(CountersSize);
+ }
+
+ if (DataSize > 0 && DevDataBegin == CachedDevDataBegin &&
+ DataSize == CachedDataSize) {
+ HostDataBegin = CachedHostData;
+ DataReused = 1;
+ if (IsVerboseMode())
+ PROF_NOTE("Reusing cached data section (%zu bytes)\n", DataSize);
+ } else if (DataSize > 0) {
+ HostDataBegin = (char *)malloc(DataSize);
+ }
+
+ if (NamesSize > 0 && DevNamesBegin == CachedDevNamesBegin &&
+ NamesSize == CachedNamesSize) {
+ HostNamesBegin = CachedHostNames;
+ NamesReused = 1;
+ if (IsVerboseMode())
+ PROF_NOTE("Reusing cached names section (%zu bytes)\n", NamesSize);
+ } else if (NamesSize > 0) {
+ HostNamesBegin = (char *)malloc(NamesSize);
+ }
+
+ // On failure before the contiguous buffer exists, free host copies and
+ // return. Do not use goto cleanup: later locals make that ill-formed C++.
+ if ((DataSize > 0 && !HostDataBegin) ||
+ (CountersSize > 0 && !HostCountersBegin) ||
+ (NamesSize > 0 && !HostNamesBegin)) {
+ PROF_ERR("%s\n", "failed to allocate host memory for device sections");
+ freeCopiedHostSections(CntsReused ? nullptr : HostCountersBegin,
+ DataReused ? nullptr : HostDataBegin,
+ NamesReused ? nullptr : HostNamesBegin);
+ return -1;
+ }
+
+ CopiedHostSectionsCleanup HostCopies(HostCountersBegin, HostDataBegin,
+ HostNamesBegin, CntsReused, DataReused,
+ NamesReused);
+
+ if ((DataSize > 0 && !DataReused &&
+ memcpyDeviceToHost(HostDataBegin, DevDataBegin, DataSize) != 0) ||
+ (CountersSize > 0 && !CntsReused &&
+ memcpyDeviceToHost(HostCountersBegin, DevCntsBegin, CountersSize) !=
+ 0) ||
+ (NamesSize > 0 && !NamesReused &&
+ memcpyDeviceToHost(HostNamesBegin, DevNamesBegin, NamesSize) != 0)) {
+ PROF_ERR("%s\n", "failed to copy profile sections from device");
+ return -1;
+ }
+
+ if (!CntsReused && CountersSize > 0) {
+ CachedDevCntsBegin = DevCntsBegin;
+ CachedHostCnts = HostCountersBegin;
+ CachedCntsSize = CountersSize;
+ }
+ if (!DataReused && DataSize > 0) {
+ CachedDevDataBegin = DevDataBegin;
+ CachedHostData = HostDataBegin;
+ CachedDataSize = DataSize;
+ }
+ if (!NamesReused && NamesSize > 0) {
+ CachedDevNamesBegin = DevNamesBegin;
+ CachedHostNames = HostNamesBegin;
+ CachedNamesSize = NamesSize;
+ }
+
+ if (IsVerboseMode())
+ PROF_NOTE("Copied device sections: Counters=%zu, Data=%zu, Names=%zu\n",
+ CountersSize, DataSize, NamesSize);
+
+ // Arrange buffer as [Data][Padding][Counters][Names] to match the layout
+ // expected by lprofWriteDataImpl (CountersDelta = CountersBegin - DataBegin).
+ const uint64_t NumData = DataSize / sizeof(__llvm_profile_data);
+ const uint64_t NumBitmapBytes = 0;
+ const uint64_t VTableSectionSize = 0;
+ const uint64_t VNamesSize = 0;
+ uint64_t PaddingBytesBeforeCounters, PaddingBytesAfterCounters,
+ PaddingBytesAfterBitmapBytes, PaddingBytesAfterNames,
+ PaddingBytesAfterVTable, PaddingBytesAfterVNames;
+
+ if (__llvm_profile_get_padding_sizes_for_counters(
+ DataSize, CountersSize, NumBitmapBytes, NamesSize, VTableSectionSize,
+ VNamesSize, &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters,
+ &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterNames,
+ &PaddingBytesAfterVTable, &PaddingBytesAfterVNames) != 0) {
+ PROF_ERR("%s\n", "failed to get padding sizes");
+ return -1;
+ }
+
+ size_t ContiguousBufferSize =
+ DataSize + PaddingBytesBeforeCounters + CountersSize + NamesSize;
+ MallocBufferCleanup ContiguousBuf(malloc(ContiguousBufferSize));
+ if (!ContiguousBuf.get()) {
+ PROF_ERR("%s\n", "failed to allocate contiguous buffer");
+ return -1;
+ }
+ char *ContiguousBuffer = ContiguousBuf.get();
+ memset(ContiguousBuffer, 0, ContiguousBufferSize);
+
+ char *BufDataBegin = ContiguousBuffer;
+ char *BufCountersBegin =
+ ContiguousBuffer + DataSize + PaddingBytesBeforeCounters;
+ char *BufNamesBegin = BufCountersBegin + CountersSize;
+
+ memcpy(BufDataBegin, HostDataBegin, DataSize);
+ memcpy(BufCountersBegin, HostCountersBegin, CountersSize);
+ memcpy(BufNamesBegin, HostNamesBegin, NamesSize);
+
+ // Relocate CounterPtr in data records for file layout.
+ // CounterPtr is device-relative offset; adjust for file layout where
+ // Data section comes first, then Counters section.
+ __llvm_profile_data *RelocatedData = (__llvm_profile_data *)BufDataBegin;
+ for (uint64_t i = 0; i < NumData; ++i) {
+ if (RelocatedData[i].CounterPtr) {
+ ptrdiff_t DeviceCounterPtrOffset = (ptrdiff_t)RelocatedData[i].CounterPtr;
+ const char *DeviceDataStructAddr =
+ (const char *)DevDataBegin + (i * sizeof(__llvm_profile_data));
+ const char *DeviceCountersAddr =
+ DeviceDataStructAddr + DeviceCounterPtrOffset;
+ ptrdiff_t OffsetIntoCountersSection =
+ DeviceCountersAddr - (const char *)DevCntsBegin;
+
+ ptrdiff_t NewRelativeOffset = DataSize + PaddingBytesBeforeCounters +
+ OffsetIntoCountersSection -
+ (i * sizeof(__llvm_profile_data));
+ memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, CounterPtr),
+ &NewRelativeOffset, sizeof(NewRelativeOffset));
+ }
+ memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, BitmapPtr),
+ 0,
+ sizeof(RelocatedData[i].BitmapPtr) +
+ sizeof(RelocatedData[i].FunctionPointer) +
+ sizeof(RelocatedData[i].Values));
+ }
+
+ char TUIndexStr[16];
+ snprintf(TUIndexStr, sizeof(TUIndexStr), "%d", TUIndex);
+
+ ret = __llvm_write_custom_profile(
+ Target, (__llvm_profile_data *)BufDataBegin,
+ (__llvm_profile_data *)(BufDataBegin + DataSize), BufCountersBegin,
+ BufCountersBegin + CountersSize, BufNamesBegin, BufNamesBegin + NamesSize,
+ nullptr);
+
+ if (ret != 0) {
+ PROF_ERR("%s\n", "failed to write device profile using shared API");
+ } else if (IsVerboseMode()) {
+ PROF_NOTE("%s\n", "Successfully wrote device profile using shared API");
+ }
+
+ return ret;
+}
+
+static int ProcessShadowVariable(void *ShadowVar, int TUIndex,
+ const char *Target) {
+ void *DevicePtrVar = nullptr;
+ if (hipGetSymbolAddress(&DevicePtrVar, ShadowVar) != 0) {
+ PROF_WARN("failed to get symbol address for shadow variable %p\n",
+ ShadowVar);
+ return -1;
+ }
+ // The shadow variable is a pointer to __llvm_profile_sections (defined
+ // in the GPU profile runtime). Dereference to get the struct address.
+ void *DeviceOffloadPrf = nullptr;
+ if (hipMemcpy(&DeviceOffloadPrf, DevicePtrVar, sizeof(void *), 2 /*DToH*/) !=
+ 0) {
+ PROF_WARN("failed to read sections pointer from shadow variable %p\n",
+ ShadowVar);
+ return -1;
+ }
+ return ProcessDeviceOffloadPrf(DeviceOffloadPrf, TUIndex, Target);
+}
+
+/* Check if HIP runtime is available and loaded */
+static int IsHipAvailable(void) {
+ EnsureHipLoaded();
+ return pHipMemcpy != nullptr && pHipGetSymbolAddress != nullptr;
+}
+
+/* -------------------------------------------------------------------------- */
+/* Collect device-side profile data */
+/* -------------------------------------------------------------------------- */
+
+extern "C" int __llvm_profile_hip_collect_device_data(void) {
+ if (NumShadowVariables == 0 && NumDynamicModules == 0)
+ return 0;
+
+ if (!IsHipAvailable())
+ return 0;
+
+ int Ret = 0;
+
+ /* Shadow variables (static-linked kernels).
+ * Iterate over all devices to collect profile data from each GPU. */
+ if (NumShadowVariables > 0) {
+ int OrigDevice = -1;
+ hipGetDevice(&OrigDevice);
+
+ for (int Dev = 0; Dev < NumDevices; ++Dev) {
+ if (hipSetDevice(Dev) != 0) {
+ if (IsVerboseMode())
+ PROF_NOTE("Failed to set device %d, skipping\n", Dev);
+ continue;
+ }
+ const char *ArchName = getDeviceArchName(Dev);
+ if (IsVerboseMode())
+ PROF_NOTE("Collecting static profile data from device %d (%s)\n", Dev,
+ ArchName);
+ for (int i = 0; i < NumShadowVariables; ++i) {
+ if (ProcessShadowVariable(OffloadShadowVariables[i], i, ArchName) != 0)
+ Ret = -1;
+ }
+ }
+
+ if (OrigDevice >= 0)
+ hipSetDevice(OrigDevice);
+ }
+
+ /* Dynamically-loaded modules — warn about any unprocessed TUs */
+ for (int i = 0; i < NumDynamicModules; ++i) {
+ OffloadDynamicModuleInfo *MI = &DynamicModules[i];
+ for (int t = 0; t < MI->NumTUs; ++t) {
+ if (!MI->TUs[t].Processed) {
+ PROF_WARN("dynamic module %p TU %d was not processed before exit\n",
+ MI->ModulePtr, t);
+ Ret = -1;
+ }
+ }
+ }
+
+ return Ret;
+}
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index dabd495cddd49..f930b649df1b1 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -23,6 +23,7 @@
#include "llvm/Analysis/CFG.h"
#include "llvm/Analysis/LoopInfo.h"
#include "llvm/Analysis/TargetLibraryInfo.h"
+#include "llvm/Frontend/Offloading/Utility.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/BasicBlock.h"
#include "llvm/IR/CFG.h"
@@ -33,12 +34,15 @@
#include "llvm/IR/DiagnosticInfo.h"
#include "llvm/IR/Dominators.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/GlobalAlias.h"
#include "llvm/IR/GlobalValue.h"
#include "llvm/IR/GlobalVariable.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instruction.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/Intrinsics.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/Type.h"
@@ -241,6 +245,20 @@ static bool profDataReferencedByCode(const Module &M) {
return enablesValueProfiling(M);
}
+// Extract CUID (Compilation Unit ID) from the module.
+// HIP/CUDA modules have a global variable __hip_cuid_<hash> that uniquely
+// identifies each translation unit. Returns empty string if not found.
+static std::string getCUIDFromModule(const Module &M) {
+ for (const GlobalVariable &GV : M.globals()) {
+ if (!GV.hasExternalLinkage())
+ continue;
+ StringRef Name = GV.getName();
+ if (Name.consume_front("__hip_cuid_"))
+ return Name.str();
+ }
+ return "";
+}
+
class InstrLowerer final {
public:
InstrLowerer(Module &M, const InstrProfOptions &Options,
@@ -265,7 +283,7 @@ class InstrLowerer final {
struct PerFunctionProfileData {
uint32_t NumValueSites[IPVK_Last + 1] = {};
GlobalVariable *RegionCounters = nullptr;
- GlobalVariable *DataVar = nullptr;
+ GlobalValue *DataVar = nullptr;
GlobalVariable *RegionBitmaps = nullptr;
uint32_t NumBitmapBytes = 0;
@@ -287,6 +305,9 @@ class InstrLowerer final {
GlobalVariable *NamesVar = nullptr;
size_t NamesSize = 0;
+ StructType *ProfileDataTy = nullptr;
+ std::string CachedCUID; // CUID cached for consistent section naming
+
// vector of counter load/store pairs to be register promoted.
std::vector<LoadStorePair> PromotionCandidates;
@@ -324,6 +345,9 @@ class InstrLowerer final {
/// Replace instrprof.increment with an increment of the appropriate value.
void lowerIncrement(InstrProfIncrementInst *Inc);
+ /// AMDGPU specific implementation of lowerIncrement.
+ void lowerIncrementAMDGPU(InstrProfIncrementInst *Inc);
+
/// Force emitting of name vars for unused functions.
void lowerCoverageData(GlobalVariable *CoverageNamesVar);
@@ -407,6 +431,25 @@ class InstrLowerer final {
/// Create a static initializer for our data, on platforms that need it,
/// and for any profile output file that was specified.
void emitInitialization();
+
+ /// For GPU targets: cache the CUID for consistent section naming.
+ void cacheGPUCUID();
+
+ /// Return the __llvm_profile_data struct type.
+ StructType *getProfileDataTy();
+
+ /// Create __llvm_offload_prf structure for GPU targets.
+ /// All sections use linker-defined __start_/__stop_ bounds.
+ void createProfileSectionSymbols();
+
+ /// Create HIP device variable registration for profile symbols
+ void createHIPDeviceVariableRegistration();
+
+ /// Create HIP dynamic module registration call
+ void createHIPDynamicModuleRegistration();
+
+ /// Create HIP dynamic module unregistration call
+ void createHIPDynamicModuleUnregistration();
};
///
@@ -938,6 +981,8 @@ bool InstrLowerer::lower() {
if (!ContainsProfiling && !CoverageNamesVar)
return MadeChange;
+ cacheGPUCUID();
+
// We did not know how many value sites there would be inside
// the instrumented function. This is counting the number of instrumented
// target value sites to enter it as field in the profile data variable.
@@ -986,6 +1031,16 @@ bool InstrLowerer::lower() {
emitNameData();
emitVTableNames();
+ // Create start/stop symbols for device code profile sections
+ createProfileSectionSymbols();
+
+ // Create host shadow variables and registration calls for HIP device profile
+ // symbols
+ createHIPDeviceVariableRegistration();
+
+ createHIPDynamicModuleRegistration();
+ createHIPDynamicModuleUnregistration();
+
// Emit runtime hook for the cases where the target does not unconditionally
// require pulling in profile runtime, and coverage is enabled on code that is
// not eliminated by the front-end, e.g. unused functions with internal
@@ -1045,7 +1100,7 @@ void InstrLowerer::lowerValueProfileInst(InstrProfValueProfileInst *Ind) {
assert(It != ProfileDataMap.end() && It->second.DataVar &&
"value profiling detected in function with no counter increment");
- GlobalVariable *DataVar = It->second.DataVar;
+ GlobalValue *DataVar = It->second.DataVar;
uint64_t ValueKind = Ind->getValueKind()->getZExtValue();
uint64_t Index = Ind->getIndex()->getZExtValue();
for (uint32_t Kind = IPVK_First; Kind < ValueKind; ++Kind)
@@ -1057,7 +1112,7 @@ void InstrLowerer::lowerValueProfileInst(InstrProfValueProfileInst *Ind) {
CallInst *Call = nullptr;
auto *TLI = &GetTLI(*Ind->getFunction());
auto *NormalizedDataVarPtr = ConstantExpr::getPointerBitCastOrAddrSpaceCast(
- DataVar, PointerType::get(M.getContext(), 0));
+ cast<Constant>(DataVar), PointerType::get(M.getContext(), 0));
// To support value profiling calls within Windows exception handlers, funclet
// information contained within operand bundles needs to be copied over to
@@ -1107,6 +1162,8 @@ GlobalVariable *InstrLowerer::getOrCreateBiasVar(StringRef VarName) {
}
Value *InstrLowerer::getCounterAddress(InstrProfCntrInstBase *I) {
+ // Note: For AMDGPU targets, lowerIncrementAMDGPU handles counter addressing
+ // directly. This function is called for non-AMDGPU targets.
auto *Counters = getOrCreateRegionCounters(I);
IRBuilder<> Builder(I);
@@ -1189,6 +1246,10 @@ void InstrLowerer::lowerTimestamp(
}
void InstrLowerer::lowerIncrement(InstrProfIncrementInst *Inc) {
+ if (TT.isAMDGPU()) {
+ lowerIncrementAMDGPU(Inc);
+ return;
+ }
auto *Addr = getCounterAddress(Inc);
IRBuilder<> Builder(Inc);
@@ -1218,6 +1279,35 @@ void InstrLowerer::lowerIncrement(InstrProfIncrementInst *Inc) {
Inc->eraseFromParent();
}
+void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
+ IRBuilder<> Builder(Inc);
+ LLVMContext &Context = M.getContext();
+ auto *Int64Ty = Type::getInt64Ty(Context);
+
+ auto *CounterIdx = Inc->getIndex();
+
+ // --- Counter address ---
+ GlobalVariable *Counters = getOrCreateRegionCounters(Inc);
+ Value *Indices[] = {Builder.getInt32(0), CounterIdx};
+ Value *Addr = Builder.CreateInBoundsGEP(Counters->getValueType(), Counters,
+ Indices, "ctr.addr");
+
+ auto *PtrTy = PointerType::getUnqual(Context);
+ Value *UniformAddrArg = ConstantPointerNull::get(cast<PointerType>(PtrTy));
+ Value *CastAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PtrTy);
+
+ Value *IncStep = Inc->getStep();
+ Value *StepI64 = Builder.CreateZExtOrTrunc(IncStep, Int64Ty, "step.i64");
+
+ auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
+ {PtrTy, PtrTy, Int64Ty}, false);
+ FunctionCallee IncrFn =
+ M.getOrInsertFunction("__llvm_profile_instrument_gpu", CalleeTy);
+ Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
+
+ Inc->eraseFromParent();
+}
+
void InstrLowerer::lowerCoverageData(GlobalVariable *CoverageNamesVar) {
ConstantArray *Names =
cast<ConstantArray>(CoverageNamesVar->getInitializer());
@@ -1400,6 +1490,12 @@ static inline Constant *getFuncAddrForProfData(Function *Fn) {
if (shouldUsePublicSymbol(Fn))
return Fn;
+ // For GPU targets, weak functions cannot use private aliases because
+ // LTO may pick a different TU's copy, leaving the alias undefined
+ if (isGPUProfTarget(*Fn->getParent()) &&
+ GlobalValue::isWeakForLinker(Fn->getLinkage()))
+ return Fn;
+
// When possible use a private alias to avoid symbolic relocations.
auto *GA = GlobalAlias::create(GlobalValue::LinkageTypes::PrivateLinkage,
Fn->getName() + ".local", Fn);
@@ -1623,11 +1719,15 @@ GlobalVariable *InstrLowerer::setupProfileSection(InstrProfInstBase *Inc,
}
Ptr->setVisibility(Visibility);
- // Put the counters and bitmaps in their own sections so linkers can
- // remove unneeded sections.
Ptr->setSection(getInstrProfSectionName(IPSK, TT.getObjectFormat()));
Ptr->setLinkage(Linkage);
- maybeSetComdat(Ptr, Fn, VarName);
+ if (isGPUProfTarget(M) && !Ptr->hasComdat()) {
+ Ptr->setComdat(M.getOrInsertComdat(VarName));
+ Ptr->setLinkage(GlobalValue::LinkOnceODRLinkage);
+ Ptr->setVisibility(GlobalValue::ProtectedVisibility);
+ } else {
+ maybeSetComdat(Ptr, Fn, VarName);
+ }
return Ptr;
}
@@ -1799,7 +1899,8 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
}
uint64_t NumCounters = Inc->getNumCounters()->getZExtValue();
- auto *CounterPtr = PD.RegionCounters;
+
+ Constant *CounterPtr = PD.RegionCounters;
uint64_t NumBitmapBytes = PD.NumBitmapBytes;
@@ -1807,11 +1908,7 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
auto *IntPtrTy = M.getDataLayout().getIntPtrType(M.getContext());
auto *Int16Ty = Type::getInt16Ty(Ctx);
auto *Int16ArrayTy = ArrayType::get(Int16Ty, IPVK_Last + 1);
- Type *DataTypes[] = {
-#define INSTR_PROF_DATA(Type, LLVMType, Name, Init) LLVMType,
-#include "llvm/ProfileData/InstrProfData.inc"
- };
- auto *DataTy = StructType::get(Ctx, ArrayRef(DataTypes));
+ auto *DataTy = getProfileDataTy();
Constant *FunctionAddr = getFuncAddrForProfData(Fn);
@@ -1819,6 +1916,15 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
for (uint32_t Kind = IPVK_First; Kind <= IPVK_Last; ++Kind)
Int16ArrayVals[Kind] = ConstantInt::get(Int16Ty, PD.NumValueSites[Kind]);
+ if (isGPUProfTarget(M)) {
+ // For GPU targets, weak functions need weak linkage for their profile data
+ // aliases to allow linker deduplication across TUs
+ if (GlobalValue::isWeakForLinker(Fn->getLinkage()))
+ Linkage = Fn->getLinkage();
+ else
+ Linkage = GlobalValue::ExternalLinkage;
+ Visibility = GlobalValue::ProtectedVisibility;
+ }
// If the data variable is not referenced by code (if we don't emit
// @llvm.instrprof.value.profile, NS will be 0), and the counter keeps the
// data variable live under linker GC, the data variable can be private. This
@@ -1830,7 +1936,8 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
// If profd is in a deduplicate comdat, NS==0 with a hash suffix guarantees
// that other copies must have the same CFG and cannot have value profiling.
// If no hash suffix, other profd copies may be referenced by code.
- if (NS == 0 && !(DataReferencedByCode && NeedComdat && !Renamed) &&
+ if (!isGPUProfTarget(M) && NS == 0 &&
+ !(DataReferencedByCode && NeedComdat && !Renamed) &&
(TT.isOSBinFormatELF() ||
(!DataReferencedByCode && TT.isOSBinFormatCOFF()))) {
Linkage = GlobalValue::PrivateLinkage;
@@ -1843,6 +1950,9 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
Visibility = GlobalValue::ProtectedVisibility;
auto *Data =
new GlobalVariable(M, DataTy, false, Linkage, nullptr, DataVarName);
+ GlobalValue *DataVar = Data;
+ Constant *DataAddr = Data;
+
Constant *RelativeCounterPtr;
GlobalVariable *BitmapPtr = PD.RegionBitmaps;
Constant *RelativeBitmapPtr = ConstantInt::get(IntPtrTy, 0);
@@ -1855,9 +1965,6 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
if (BitmapPtr != nullptr)
RelativeBitmapPtr = ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy);
} else if (TT.isNVPTX()) {
- // The NVPTX target cannot handle self-referencing constant expressions in
- // global initializers at all. Use absolute pointers and have the runtime
- // registration convert them to relative offsets.
DataSectionKind = IPSK_data;
RelativeCounterPtr = ConstantExpr::getPtrToInt(CounterPtr, IntPtrTy);
} else {
@@ -1866,29 +1973,36 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
DataSectionKind = IPSK_data;
RelativeCounterPtr =
ConstantExpr::getSub(ConstantExpr::getPtrToInt(CounterPtr, IntPtrTy),
- ConstantExpr::getPtrToInt(Data, IntPtrTy));
+ ConstantExpr::getPtrToInt(DataAddr, IntPtrTy));
if (BitmapPtr != nullptr)
RelativeBitmapPtr =
ConstantExpr::getSub(ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy),
- ConstantExpr::getPtrToInt(Data, IntPtrTy));
+ ConstantExpr::getPtrToInt(DataAddr, IntPtrTy));
}
Constant *DataVals[] = {
#define INSTR_PROF_DATA(Type, LLVMType, Name, Init) Init,
#include "llvm/ProfileData/InstrProfData.inc"
};
- Data->setInitializer(ConstantStruct::get(DataTy, DataVals));
+ auto *DataInit = ConstantStruct::get(DataTy, DataVals);
- Data->setVisibility(Visibility);
- Data->setSection(
+ auto *DataGV = cast<GlobalVariable>(DataVar);
+ DataGV->setInitializer(DataInit);
+ DataGV->setVisibility(Visibility);
+ DataGV->setSection(
getInstrProfSectionName(DataSectionKind, TT.getObjectFormat()));
- Data->setAlignment(Align(INSTR_PROF_DATA_ALIGNMENT));
- maybeSetComdat(Data, Fn, CntsVarName);
+ DataGV->setAlignment(Align(INSTR_PROF_DATA_ALIGNMENT));
+ if (isGPUProfTarget(M) && !DataGV->hasComdat()) {
+ DataGV->setComdat(M.getOrInsertComdat(CntsVarName));
+ DataGV->setLinkage(GlobalValue::LinkOnceODRLinkage);
+ } else {
+ maybeSetComdat(DataGV, Fn, CntsVarName);
+ }
- PD.DataVar = Data;
+ PD.DataVar = DataVar;
// Mark the data variable as used so that it isn't stripped out.
- CompilerUsedVars.push_back(Data);
+ CompilerUsedVars.push_back(DataVar);
// Now that the linkage set by the FE has been passed to the data and counter
// variables, reset Name variable's linkage and visibility to private so that
// it can be removed later by the compiler.
@@ -1948,6 +2062,102 @@ void InstrLowerer::emitVNodes() {
UsedVars.push_back(VNodesVar);
}
+void InstrLowerer::createHIPDynamicModuleRegistration() {
+ if (isGPUProfTarget(M))
+ return;
+ StringRef FuncNames[] = {"hipModuleLoad", "hipModuleLoadData",
+ "hipModuleLoadDataEx"};
+ for (StringRef FuncName : FuncNames) {
+ Function *F = M.getFunction(FuncName);
+ if (!F)
+ continue;
+
+ for (User *U : F->users()) {
+ if (auto *CB = dyn_cast<CallBase>(U)) {
+ Instruction *InsertPt = nullptr;
+ // If the call is an invoke instruction, we should insert the
+ // registration call in the normal destination block.
+ if (auto *Invoke = dyn_cast<InvokeInst>(CB)) {
+ InsertPt = &*Invoke->getNormalDest()->getFirstInsertionPt();
+ } else if (CB->isTerminator()) {
+ // If it's another kind of terminator (e.g., callbr), we don't
+ // know the semantics of the successors, so we conservatively
+ // skip it. The hipModuleLoad* functions are not expected to be
+ // used in other terminator instructions.
+ continue;
+ } else {
+ // This is a normal call instruction, so we can insert after it.
+ InsertPt = CB->getNextNode();
+ }
+
+ // If there's no valid insertion point (e.g., a malformed block),
+ // skip.
+ if (!InsertPt)
+ continue;
+
+ IRBuilder<> Builder(InsertPt);
+ auto *VoidTy = Type::getVoidTy(M.getContext());
+ auto *VoidPtrTy = PointerType::getUnqual(M.getContext());
+ auto *Int32Ty = Type::getInt32Ty(M.getContext());
+ // register(int rc, void **modulePtr, const void *image)
+ auto *RegisterDynamicModuleTy =
+ FunctionType::get(VoidTy, {Int32Ty, VoidPtrTy, VoidPtrTy}, false);
+ FunctionCallee RegisterFunc = M.getOrInsertFunction(
+ "__llvm_profile_offload_register_dynamic_module",
+ RegisterDynamicModuleTy);
+
+ // Arg 0: return value of the hipModuleLoad* call (hipError_t / i32).
+ Value *ReturnValue = CB;
+ // Arg 1: module handle (out-parameter, hipModule_t*).
+ Value *ModuleHandle = CB->getArgOperand(0);
+ // Arg 2: code object image pointer.
+ // For hipModuleLoadData(module, image) and
+ // hipModuleLoadDataEx(module, image, ...), image is arg 1.
+ // For hipModuleLoad(module, fname), arg 1 is a filename — pass NULL.
+ Value *ImagePtr;
+ if (FuncName == "hipModuleLoad")
+ ImagePtr =
+ ConstantPointerNull::get(PointerType::getUnqual(M.getContext()));
+ else
+ ImagePtr = CB->getArgOperand(1);
+
+ Builder.CreateCall(RegisterFunc, {ReturnValue, ModuleHandle, ImagePtr});
+ }
+ }
+ }
+}
+
+void InstrLowerer::createHIPDynamicModuleUnregistration() {
+ Function *F = M.getFunction("hipModuleUnload");
+ if (!F)
+ return;
+
+ for (User *U : F->users()) {
+ if (auto *CB = dyn_cast_or_null<CallBase>(U)) {
+ // The insertion point is right before the call to hipModuleUnload.
+ Instruction *InsertPt = CB;
+
+ IRBuilder<> Builder(InsertPt);
+ auto *VoidTy = Type::getVoidTy(M.getContext());
+ auto *VoidPtrTy = PointerType::getUnqual(M.getContext());
+
+ auto *UnregisterDynamicModuleTy =
+ FunctionType::get(VoidTy, {VoidPtrTy}, false);
+ FunctionCallee UnregisterFunc = M.getOrInsertFunction(
+ "__llvm_profile_offload_unregister_dynamic_module",
+ UnregisterDynamicModuleTy);
+
+ // The argument is the module handle, which is the first
+ // argument to the hipModuleUnload call.
+ Value *ModuleHandle = CB->getArgOperand(0);
+ Value *CastedModuleHandle =
+ Builder.CreatePointerCast(ModuleHandle, VoidPtrTy);
+
+ Builder.CreateCall(UnregisterFunc, {CastedModuleHandle});
+ }
+ }
+}
+
void InstrLowerer::emitNameData() {
if (ReferencedNames.empty())
return;
@@ -1961,16 +2171,23 @@ void InstrLowerer::emitNameData() {
auto &Ctx = M.getContext();
auto *NamesVal =
ConstantDataArray::getString(Ctx, StringRef(CompressedNameStr), false);
- NamesVar = new GlobalVariable(M, NamesVal->getType(), true,
- GlobalValue::PrivateLinkage, NamesVal,
- getInstrProfNamesVarName());
+ std::string NamesVarName = std::string(getInstrProfNamesVarName());
+ if (isGPUProfTarget(M)) {
+ std::string CUID = CachedCUID.empty() ? getCUIDFromModule(M) : CachedCUID;
+ if (!CUID.empty())
+ NamesVarName = NamesVarName + "_" + CUID;
+ }
+ NamesVar =
+ new GlobalVariable(M, NamesVal->getType(), true,
+ GlobalValue::PrivateLinkage, NamesVal, NamesVarName);
NamesSize = CompressedNameStr.size();
setGlobalVariableLargeSection(TT, *NamesVar);
- NamesVar->setSection(
+ std::string NamesSectionName =
ProfileCorrelate == InstrProfCorrelator::BINARY
? getInstrProfSectionName(IPSK_covname, TT.getObjectFormat())
- : getInstrProfSectionName(IPSK_name, TT.getObjectFormat()));
+ : getInstrProfSectionName(IPSK_name, TT.getObjectFormat());
+ NamesVar->setSection(NamesSectionName);
// On COFF, it's important to reduce the alignment down to 1 to prevent the
// linker from inserting padding before the start of the names section or
// between names entries.
@@ -2179,3 +2396,188 @@ void createProfileSamplingVar(Module &M) {
appendToCompilerUsed(M, SamplingVar);
}
} // namespace llvm
+
+namespace {
+
+// For GPU targets: Allocate contiguous arrays for all profile data.
+// This solves the linker reordering problem by using ONE symbol per section
+// type, so there's nothing for the linker to reorder.
+StructType *InstrLowerer::getProfileDataTy() {
+ if (ProfileDataTy)
+ return ProfileDataTy;
+
+ auto &Ctx = M.getContext();
+ auto *IntPtrTy = M.getDataLayout().getIntPtrType(M.getContext());
+ auto *Int16Ty = Type::getInt16Ty(Ctx);
+ auto *Int16ArrayTy = ArrayType::get(Int16Ty, IPVK_Last + 1);
+ Type *DataTypes[] = {
+#define INSTR_PROF_DATA(Type, LLVMType, Name, Init) LLVMType,
+#include "llvm/ProfileData/InstrProfData.inc"
+ };
+ ProfileDataTy = StructType::get(Ctx, ArrayRef(DataTypes));
+ return ProfileDataTy;
+}
+
+void InstrLowerer::cacheGPUCUID() {
+ if (!isGPUProfTarget(M))
+ return;
+ CachedCUID = getCUIDFromModule(M);
+}
+
+// Create CUID-suffixed pointer to __llvm_profile_sections for GPU targets.
+// The basic HIP/offload runtime exposes the original 7-entry section table:
+// names_start, names_stop, cnts_start, cnts_stop, data_start, data_stop,
+// raw_version.
+// We create a per-TU global that points to it, giving the host a unique
+// symbol for shadow variable registration.
+void InstrLowerer::createProfileSectionSymbols() {
+ if (!isGPUProfTarget(M) || CachedCUID.empty())
+ return;
+
+ auto &Ctx = M.getContext();
+ unsigned AS = M.getDataLayout().getDefaultGlobalsAddressSpace();
+ auto *Int8PtrTy = PointerType::get(Ctx, AS);
+
+ // __llvm_profile_sections is an array of 7 pointers defined in the GPU
+ // profile runtime (InstrProfilingPlatformGPU.c). Declare it as external.
+ auto *SectionsTy = ArrayType::get(Int8PtrTy, 7);
+ auto *SectionsGV = M.getGlobalVariable("__llvm_profile_sections");
+ if (!SectionsGV) {
+ SectionsGV = new GlobalVariable(M, SectionsTy, /*isConstant=*/true,
+ GlobalValue::ExternalLinkage, nullptr,
+ "__llvm_profile_sections", nullptr,
+ GlobalValue::NotThreadLocal, AS);
+ SectionsGV->setVisibility(GlobalValue::HiddenVisibility);
+ }
+
+ // Create a CUID-suffixed global that stores a pointer to the sections
+ // struct. Aliases can't point to declarations, so we use a pointer global.
+ // The host reads through this indirection: hipGetSymbolAddress gives the
+ // pointer global's device address, then one DtoH copy yields the sections
+ // struct address, then another DtoH copy reads the actual sections.
+ auto *PtrTy = PointerType::get(Ctx, AS);
+ auto *PtrInit =
+ ConstantExpr::getPointerBitCastOrAddrSpaceCast(SectionsGV, PtrTy);
+ std::string PtrName = "__llvm_offload_prf_" + CachedCUID;
+ auto *PtrGV = new GlobalVariable(
+ M, PtrTy, /*isConstant=*/true, GlobalValue::ExternalLinkage, PtrInit,
+ PtrName, nullptr, GlobalValue::NotThreadLocal, AS);
+ PtrGV->setVisibility(GlobalValue::DefaultVisibility);
+ CompilerUsedVars.push_back(PtrGV);
+}
+
+void InstrLowerer::createHIPDeviceVariableRegistration() {
+ if (isGPUProfTarget(M))
+ return;
+
+ std::string CUID = CachedCUID.empty() ? getCUIDFromModule(M) : CachedCUID;
+ if (CUID.empty())
+ return;
+
+ auto &Ctx = M.getContext();
+ auto *VoidTy = Type::getVoidTy(Ctx);
+ auto *VoidPtrTy = PointerType::getUnqual(Ctx);
+
+ std::string OffloadPrfName = "__llvm_offload_prf_" + CUID;
+ auto *OffloadPrfShadow = new GlobalVariable(
+ M, VoidPtrTy, /*isConstant=*/false, GlobalValue::ExternalLinkage,
+ ConstantPointerNull::get(cast<PointerType>(VoidPtrTy)), OffloadPrfName);
+ CompilerUsedVars.push_back(OffloadPrfShadow);
+
+ auto *RegisterShadowTy = FunctionType::get(VoidTy, {VoidPtrTy}, false);
+ FunctionCallee RegisterShadowFunc = M.getOrInsertFunction(
+ "__llvm_profile_offload_register_shadow_variable", RegisterShadowTy);
+
+ Function *Ctor = M.getFunction("__hip_module_ctor");
+ if (!Ctor) {
+ // RDC mode: no __hip_module_ctor per-TU. Emit an offloading entry so the
+ // linker wrapper generates __hipRegisterVar in the final module ctor.
+ llvm::offloading::emitOffloadingEntry(
+ M, llvm::object::OffloadKind::OFK_HIP, OffloadPrfShadow, OffloadPrfName,
+ M.getDataLayout().getPointerSize(),
+ llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
+
+ auto *CtorFn = Function::Create(FunctionType::get(VoidTy, false),
+ GlobalValue::InternalLinkage,
+ "__llvm_pgo_register_" + CUID, &M);
+ auto *Entry = BasicBlock::Create(Ctx, "entry", CtorFn);
+ IRBuilder<> B(Entry);
+ B.CreateCall(RegisterShadowFunc, {OffloadPrfShadow});
+ B.CreateRetVoid();
+ appendToGlobalCtors(M, CtorFn, 65535);
+ return;
+ }
+
+ // Locate the HIP fat-binary registration call and capture its return value
+ Value *Handle = nullptr;
+ for (BasicBlock &BB : *Ctor)
+ for (Instruction &I : BB)
+ if (auto *CB = dyn_cast<CallBase>(&I))
+ if (Function *Callee = CB->getCalledFunction())
+ if (Callee->getName() == "__hipRegisterFatBinary") {
+ Handle = &I; // call result
+ break;
+ }
+ if (!Handle)
+ return;
+ GlobalVariable *FatbinHandleGV = nullptr;
+ if (auto *HandleInst = dyn_cast<Instruction>(Handle))
+ for (Instruction *Cur = HandleInst->getNextNode(); Cur;
+ Cur = Cur->getNextNode()) {
+ auto *SI = dyn_cast<StoreInst>(Cur);
+ if (!SI || SI->getValueOperand() != Handle)
+ continue;
+ if (auto *GV = dyn_cast<GlobalVariable>(
+ SI->getPointerOperand()->stripPointerCasts())) {
+ FatbinHandleGV = GV;
+ break;
+ }
+ }
+
+ if (!FatbinHandleGV) {
+ LLVM_DEBUG(llvm::dbgs()
+ << "store of __hipRegisterFatBinary call not found\n");
+ }
+
+ // Insert the new registration just before the ctor’s return
+ ReturnInst *RetInst = nullptr;
+ for (auto &BB : llvm::reverse(*Ctor))
+ if ((RetInst = dyn_cast<ReturnInst>(BB.getTerminator())))
+ break;
+ if (!RetInst)
+ return;
+ IRBuilder<> Builder(RetInst);
+
+ LLVM_DEBUG(
+ llvm::dbgs() << "Found __hip_module_ctor, registering anchors for CUID="
+ << CUID << "\n");
+
+ auto *Int32Ty = Type::getInt32Ty(Ctx);
+ auto *Int64Ty = Type::getInt64Ty(Ctx);
+ auto *RegisterVarTy =
+ FunctionType::get(VoidTy,
+ {VoidPtrTy, VoidPtrTy, VoidPtrTy, VoidPtrTy, Int32Ty,
+ Int64Ty, Int32Ty, Int32Ty},
+ false);
+ FunctionCallee RegisterVarFunc =
+ M.getOrInsertFunction("__hipRegisterVar", RegisterVarTy);
+ Value *HipHandle =
+ FatbinHandleGV ? Builder.CreateLoad(VoidPtrTy, FatbinHandleGV) : Handle;
+
+ auto *NameStr = ConstantDataArray::getString(Ctx, OffloadPrfName, true);
+ auto *NameGV = new GlobalVariable(M, NameStr->getType(), true,
+ GlobalValue::PrivateLinkage, NameStr,
+ OffloadPrfName + ".name");
+
+ Builder.CreateCall(RegisterVarFunc,
+ {HipHandle, OffloadPrfShadow,
+ Builder.CreatePointerCast(NameGV, VoidPtrTy),
+ Builder.CreatePointerCast(NameGV, VoidPtrTy),
+ Builder.getInt32(0),
+ Builder.getInt64(M.getDataLayout().getPointerSize()),
+ Builder.getInt32(0), Builder.getInt32(0)});
+
+ Builder.CreateCall(RegisterShadowFunc, {OffloadPrfShadow});
+}
+
+} // namespace
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index db032d6fcad45..404c47d5183f1 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -379,8 +379,10 @@ class FunctionInstrumenter final {
// values. Supporting other values is relatively straight-forward - just
// another counter range within the context.
bool isValueProfilingDisabled() const {
+ const Triple &TT = M.getTargetTriple();
return DisableValueProfiling ||
- InstrumentationType == PGOInstrumentationType::CTXPROF;
+ InstrumentationType == PGOInstrumentationType::CTXPROF ||
+ TT.isAMDGPU() || TT.isNVPTX();
}
bool shouldInstrumentEntryBB() const {
@@ -470,7 +472,7 @@ createIRLevelProfileFlagVar(Module &M,
Constant::getIntegerValue(IntTy64, APInt(64, ProfileVersion)), VarName);
IRLevelVersionVariable->setVisibility(GlobalValue::HiddenVisibility);
- Triple TT(M.getTargetTriple());
+ const Triple &TT = M.getTargetTriple();
if (TT.supportsCOMDAT()) {
IRLevelVersionVariable->setLinkage(GlobalValue::ExternalLinkage);
IRLevelVersionVariable->setComdat(M.getOrInsertComdat(VarName));
@@ -1959,7 +1961,7 @@ static bool InstrumentAllFunctions(
if (InstrumentationType == PGOInstrumentationType::FDO)
createIRLevelProfileFlagVar(M, InstrumentationType);
- Triple TT(M.getTargetTriple());
+ const Triple &TT = M.getTargetTriple();
LLVMContext &Ctx = M.getContext();
if (!TT.isOSBinFormatELF() && EnableVTableValueProfiling)
Ctx.diagnose(DiagnosticInfoPGOProfile(
@@ -2406,8 +2408,7 @@ void llvm::setProfMetadata(Instruction *TI, ArrayRef<uint64_t> EdgeCounts,
uint64_t MaxCount) {
auto Weights = downscaleWeights(EdgeCounts, MaxCount);
- LLVM_DEBUG(dbgs() << "Weight is: "; for (const auto &W
- : Weights) {
+ LLVM_DEBUG(dbgs() << "Weight is: "; for (const auto &W : Weights) {
dbgs() << W << " ";
} dbgs() << "\n";);
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
new file mode 100644
index 0000000000000..2252a722a91cc
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
@@ -0,0 +1,34 @@
+;; Test that AMDGPU targets use contiguous counter allocation with CUID-based naming.
+;; This avoids linker reordering issues where individual __profc_* symbols could be
+;; placed in any order within the section.
+
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -passes=instrprof < %s | FileCheck %s
+
+;; Simulate a module with CUID (as generated by HIP compilation)
+ at __hip_cuid_abc123 = addrspace(1) global i8 0
+
+ at __profn_kernel1 = private constant [7 x i8] c"kernel1"
+ at __profn_kernel2 = private constant [7 x i8] c"kernel2"
+
+;; Per-kernel counter arrays: kernel1 has 2 slots, kernel2 has 1 (section "__llvm_prf_cnts")
+; CHECK: @__profc_kernel1 = linkonce_odr protected addrspace(1) global [2 x i64] zeroinitializer, section "__llvm_prf_cnts"
+; CHECK: @__profc_kernel2 = linkonce_odr protected addrspace(1) global [1 x i64] zeroinitializer, section "__llvm_prf_cnts"
+
+define amdgpu_kernel void @kernel1() {
+ call void @llvm.instrprof.increment(ptr @__profn_kernel1, i64 12345, i32 2, i32 0)
+ call void @llvm.instrprof.increment(ptr @__profn_kernel1, i64 12345, i32 2, i32 1)
+ ret void
+}
+
+define amdgpu_kernel void @kernel2() {
+ call void @llvm.instrprof.increment(ptr @__profn_kernel2, i64 67890, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+;; Registration symbol points at external section table (Joseph #187136 layout)
+; CHECK: @__llvm_offload_prf_abc123 = addrspace(1) constant ptr addrspace(1) @__llvm_profile_sections
+
+;; Second counter slot uses GEP into the same [2 x i64] arrays
+; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([2 x i64], ptr addrspace(1) @__profc_kernel1, i32 0, i32 1) to ptr), ptr null, i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
new file mode 100644
index 0000000000000..efe53ab1ebdfb
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-instrumentation.ll
@@ -0,0 +1,32 @@
+;; Test basic AMDGPU PGO instrumentation lowering.
+;; Verifies that each instrumentation point lowers directly to a call to
+;; __llvm_profile_instrument_gpu with a null uniform-counter argument.
+
+; RUN: opt %s -mtriple=amdgcn-amd-amdhsa -passes=instrprof -S | FileCheck %s
+
+ at __hip_cuid_test01 = addrspace(1) global i8 0
+ at __profn_test_kernel = private constant [11 x i8] c"test_kernel"
+
+define amdgpu_kernel void @test_kernel(ptr addrspace(1) %out, i32 %n) {
+entry:
+ call void @llvm.instrprof.increment(ptr @__profn_test_kernel, i64 111, i32 4, i32 0)
+ %cmp = icmp sgt i32 %n, 0
+ br i1 %cmp, label %if.then, label %if.end
+
+if.then:
+ call void @llvm.instrprof.increment(ptr @__profn_test_kernel, i64 111, i32 4, i32 1)
+ store i32 1, ptr addrspace(1) %out
+ br label %if.end
+
+if.end:
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
+
+; CHECK-LABEL: define {{.*}} @test_kernel
+; CHECK-NOT: @__llvm_profile_sampling_gpu
+; CHECK: call void @__llvm_profile_instrument_gpu(
+; CHECK-SAME: ptr addrspacecast (ptr addrspace(1) @__profc_test_kernel to ptr), ptr null, i64 1)
+; CHECK: call void @__llvm_profile_instrument_gpu(
+; CHECK-SAME: ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([4 x i64], ptr addrspace(1) @__profc_test_kernel, i32 0, i32 1) to ptr), ptr null, i64 1)
diff --git a/llvm/test/Instrumentation/InstrProfiling/gpu-weak.ll b/llvm/test/Instrumentation/InstrProfiling/gpu-weak.ll
new file mode 100644
index 0000000000000..ce16f1ee3215f
--- /dev/null
+++ b/llvm/test/Instrumentation/InstrProfiling/gpu-weak.ll
@@ -0,0 +1,36 @@
+; RUN: opt < %s -passes=instrprof -S | FileCheck %s
+
+; Test that weak functions on GPU targets get weak linkage for their
+; __profd_ aliases to allow linker deduplication across TUs.
+; Non-weak functions get external linkage (default for aliases).
+
+target triple = "amdgcn-amd-amdhsa"
+
+ at __hip_cuid_abc123 = addrspace(1) global i8 0
+
+; AMDGPU GPU profiling lowers to per-function comdat globals (not aliases).
+; CHECK: @__profd_weak_func = linkonce_odr protected addrspace(1) global
+ at __profn_weak_func = private constant [9 x i8] c"weak_func"
+
+define weak void @weak_func() {
+ call void @llvm.instrprof.increment(ptr @__profn_weak_func, i64 0, i32 1, i32 0)
+ ret void
+}
+
+; CHECK: @__profd_weak_odr_func = linkonce_odr protected addrspace(1) global
+ at __profn_weak_odr_func = private constant [13 x i8] c"weak_odr_func"
+
+define weak_odr void @weak_odr_func() {
+ call void @llvm.instrprof.increment(ptr @__profn_weak_odr_func, i64 0, i32 1, i32 0)
+ ret void
+}
+
+; CHECK: @__profd_normal_func = linkonce_odr protected addrspace(1) global
+ at __profn_normal_func = private constant [11 x i8] c"normal_func"
+
+define void @normal_func() {
+ call void @llvm.instrprof.increment(ptr @__profn_normal_func, i64 0, i32 1, i32 0)
+ ret void
+}
+
+declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
diff --git a/llvm/test/Transforms/PGOProfile/amdgpu-disable-value-profiling.ll b/llvm/test/Transforms/PGOProfile/amdgpu-disable-value-profiling.ll
new file mode 100644
index 0000000000000..21b1d68004b13
--- /dev/null
+++ b/llvm/test/Transforms/PGOProfile/amdgpu-disable-value-profiling.ll
@@ -0,0 +1,22 @@
+;; Test that value profiling (indirect call profiling) is disabled for GPU targets.
+;; The device-side profiling runtime does not implement
+;; __llvm_profile_instrument_target, so indirect call profiling must not be emitted.
+
+; RUN: opt < %s -passes=pgo-instr-gen -S | FileCheck %s
+
+target triple = "amdgcn-amd-amdhsa"
+
+ at fptr = addrspace(1) global ptr null, align 8
+
+;; Verify that regular block instrumentation IS emitted
+; CHECK: call void @llvm.instrprof.increment
+
+;; Verify that value profiling for indirect calls is NOT emitted
+; CHECK-NOT: call void @llvm.instrprof.value.profile
+
+define amdgpu_kernel void @test_indirect_call() {
+entry:
+ %fp = load ptr, ptr addrspace(1) @fptr, align 8
+ call void %fp()
+ ret void
+}
>From 29f38460cda466df5bd6c0c5e90460d0127369d1 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 17 Apr 2026 11:15:20 -0400
Subject: [PATCH 02/13] [PGO][AMDGPU] Apply review feedback on
InstrProfilingPlatformROCm.cpp
Address review comments on PR #177665:
1. Rename PascalCase static helpers to lowerCamelCase to match LLVM style:
EnsureHipLoaded -> ensureHipLoaded
IsVerboseMode -> isVerboseMode
UnwrapOffloadBundle -> unwrapOffloadBundle
EnumerateElfSymbols -> enumerateElfSymbols
RegisterPrfSymbol -> registerPrfSymbol
ProcessDeviceOffloadPrf -> processDeviceOffloadPrf
GrowPtrArray -> growPtrArray
ProcessShadowVariable -> processShadowVariable
IsHipAvailable -> isHipAvailable
2. In unwrapOffloadBundle, use sizeof(var) instead of hard-coded 8 for the
three entry-table field memcpys (EntryOffset, EntrySize, IDSize) and for
the corresponding Cursor advances.
No functional change.
---
.../profile/InstrProfilingPlatformROCm.cpp | 94 +++++++++----------
1 file changed, 47 insertions(+), 47 deletions(-)
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index c0b1732c8d75b..79646f17b29e2 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -30,10 +30,10 @@ extern "C" {
#include <stdlib.h>
#include <string.h>
-static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
+static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
const char *Target);
-static int IsVerboseMode() {
+static int isVerboseMode() {
static int IsVerbose = -1;
if (IsVerbose == -1)
IsVerbose = getenv("LLVM_PROFILE_VERBOSE") != nullptr;
@@ -78,14 +78,14 @@ static char DeviceArchNames[MAX_DEVICES][256];
/* Keep HIP-only to avoid an HSA dependency. */
/* -------------------------------------------------------------------------- */
-static void EnsureHipLoaded(void) {
+static void ensureHipLoaded(void) {
static int Initialized = 0;
if (Initialized)
return;
Initialized = 1;
if (!__interception::DynamicLoaderAvailable()) {
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("%s", "Dynamic library loading not available - "
"HIP profiling disabled\n");
return;
@@ -132,7 +132,7 @@ static void EnsureHipLoaded(void) {
strncpy(DeviceArchNames[i], Prop.gcnArchName,
sizeof(DeviceArchNames[i]) - 1);
DeviceArchNames[i][sizeof(DeviceArchNames[i]) - 1] = '\0';
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Device %d arch: %s\n", i, DeviceArchNames[i]);
}
}
@@ -146,13 +146,13 @@ static void EnsureHipLoaded(void) {
/* -------------------------------------------------------------------------- */
static int hipGetSymbolAddress(void **devPtr, const void *symbol) {
- EnsureHipLoaded();
+ ensureHipLoaded();
return pHipGetSymbolAddress ? pHipGetSymbolAddress(devPtr, symbol) : -1;
}
static int hipMemcpy(void *dest, const void *src, size_t len,
int kind /*2=DToH*/) {
- EnsureHipLoaded();
+ ensureHipLoaded();
return pHipMemcpy ? pHipMemcpy(dest, src, len, kind) : -1;
}
@@ -165,18 +165,18 @@ static int memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) {
static int hipModuleGetGlobal(void **DevPtr, size_t *Bytes, void *Module,
const char *Name) {
- EnsureHipLoaded();
+ ensureHipLoaded();
return pHipModuleGetGlobal ? pHipModuleGetGlobal(DevPtr, Bytes, Module, Name)
: -1;
}
static int hipGetDevice(int *DeviceId) {
- EnsureHipLoaded();
+ ensureHipLoaded();
return pHipGetDevice ? pHipGetDevice(DeviceId) : -1;
}
static int hipSetDevice(int DeviceId) {
- EnsureHipLoaded();
+ ensureHipLoaded();
return pHipSetDevice ? pHipSetDevice(DeviceId) : -1;
}
@@ -227,7 +227,7 @@ typedef int (*SymbolCallback)(const char *Name, void *UserData);
/* If Image is a clang offload bundle (__CLANG_OFFLOAD_BUNDLE__), find the
* first embedded code object that is a valid ELF and return a pointer to it.
* Otherwise return Image unchanged. Returns nullptr if no ELF is found. */
-static const void *UnwrapOffloadBundle(const void *Image) {
+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. */
@@ -240,12 +240,12 @@ static const void *UnwrapOffloadBundle(const void *Image) {
const char *Cursor = Buf + 32;
for (uint64_t I = 0; I < NumEntries; ++I) {
uint64_t EntryOffset, EntrySize, IDSize;
- memcpy(&EntryOffset, Cursor, 8);
- Cursor += 8;
- memcpy(&EntrySize, Cursor, 8);
- Cursor += 8;
- memcpy(&IDSize, Cursor, 8);
- Cursor += 8;
+ memcpy(&EntryOffset, Cursor, sizeof(EntryOffset));
+ Cursor += sizeof(EntryOffset);
+ memcpy(&EntrySize, Cursor, sizeof(EntrySize));
+ Cursor += sizeof(EntrySize);
+ memcpy(&IDSize, Cursor, sizeof(IDSize));
+ Cursor += sizeof(IDSize);
/* Skip the entry ID string. */
Cursor += IDSize;
@@ -267,20 +267,20 @@ static const void *UnwrapOffloadBundle(const void *Image) {
* name starts with PREFIX. Image may be nullptr (e.g. hipModuleLoad from file)
* or a clang offload bundle containing an ELF;
* in that case the function unwraps the bundle first. */
-static void EnumerateElfSymbols(const void *Image, const char *Prefix,
+static void enumerateElfSymbols(const void *Image, const char *Prefix,
SymbolCallback CB, void *UserData) {
if (!Image)
return;
/* Handle clang offload bundle wrapping. */
- Image = UnwrapOffloadBundle(Image);
+ 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())
+ if (isVerboseMode())
PROF_NOTE("%s", "Image is not a valid ELF, skipping enumeration\n");
return;
}
@@ -319,7 +319,7 @@ typedef struct {
/* Grow the TU array inside a module entry and register one __llvm_offload_prf_*
* symbol. Also pre-registers the corresponding per-TU section symbols with CLR
* (needed so hipMemcpy can copy from those device addresses later). */
-static int RegisterPrfSymbol(const char *Name, void *UserData) {
+static int registerPrfSymbol(const char *Name, void *UserData) {
EnumState *S = (EnumState *)UserData;
OffloadDynamicModuleInfo *MI = S->ModInfo;
@@ -367,7 +367,7 @@ static int RegisterPrfSymbol(const char *Name, void *UserData) {
extern "C" void
__llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
const void *Image) {
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n",
NumDynamicModules, ModuleLoadRc, *Ptr, Image);
@@ -398,17 +398,17 @@ __llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
* profiling is not yet supported. */
#if __has_include(<elf.h>)
EnumState State = {*Ptr, MI};
- EnumerateElfSymbols(Image, "__llvm_offload_prf_", RegisterPrfSymbol, &State);
+ enumerateElfSymbols(Image, "__llvm_offload_prf_", registerPrfSymbol, &State);
#else
(void)Image;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("%s",
"Dynamic module profiling not supported on this platform\n");
#endif
if (MI->NumTUs == 0) {
PROF_WARN("no __llvm_offload_prf_* symbols found in module %p\n", *Ptr);
- } else if (IsVerboseMode()) {
+ } else if (isVerboseMode()) {
PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs);
}
}
@@ -420,7 +420,7 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
if (MI->ModulePtr != Ptr)
continue;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Unregistering module %p (%d TUs)\n", MI->ModulePtr,
MI->NumTUs);
@@ -428,7 +428,7 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
for (int t = 0; t < MI->NumTUs; ++t) {
OffloadDynamicTUInfo *TU = &MI->TUs[t];
if (TU->Processed) {
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Module %p TU %d already processed, skipping\n", Ptr, t);
continue;
}
@@ -438,7 +438,7 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
int CurDev = 0;
hipGetDevice(&CurDev);
const char *ArchName = getDeviceArchName(CurDev);
- if (ProcessDeviceOffloadPrf(TU->DeviceVar, TUIndex, ArchName) == 0)
+ if (processDeviceOffloadPrf(TU->DeviceVar, TUIndex, ArchName) == 0)
TU->Processed = 1;
else
PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr,
@@ -448,12 +448,12 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
return;
}
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_WARN("unregister called for unknown module %p\n", Ptr);
}
/* Grow a void* array, doubling capacity (or starting at InitCap). */
-static int GrowPtrArray(void ***Arr, int *Num, int *Cap, int InitCap) {
+static int growPtrArray(void ***Arr, int *Num, int *Cap, int InitCap) {
if (*Num < *Cap)
return 0;
int NewCap = *Cap ? *Cap * 2 : InitCap;
@@ -470,7 +470,7 @@ static int NumShadowVariables = 0;
static int CapShadowVariables = 0;
extern "C" void __llvm_profile_offload_register_shadow_variable(void *ptr) {
- if (GrowPtrArray(&OffloadShadowVariables, &NumShadowVariables,
+ if (growPtrArray(&OffloadShadowVariables, &NumShadowVariables,
&CapShadowVariables, 64))
return;
OffloadShadowVariables[NumShadowVariables++] = ptr;
@@ -482,7 +482,7 @@ static int CapSectionShadowVariables = 0;
extern "C" void
__llvm_profile_offload_register_section_shadow_variable(void *ptr) {
- if (GrowPtrArray(&OffloadSectionShadowVariables, &NumSectionShadowVariables,
+ if (growPtrArray(&OffloadSectionShadowVariables, &NumSectionShadowVariables,
&CapSectionShadowVariables, 64))
return;
OffloadSectionShadowVariables[NumSectionShadowVariables++] = ptr;
@@ -536,7 +536,7 @@ struct MallocBufferCleanup {
} // namespace
-static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
+static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
const char *Target) {
__llvm_profile_gpu_sections HostSections;
@@ -557,7 +557,7 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
size_t DataSize = (const char *)DevDataEnd - (const char *)DevDataBegin;
size_t NamesSize = (const char *)DevNamesEnd - (const char *)DevNamesBegin;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Section pointers: Cnts=[%p,%p]=%zu Data=[%p,%p]=%zu "
"Names=[%p,%p]=%zu\n",
DevCntsBegin, DevCntsEnd, CountersSize, DevDataBegin, DevDataEnd,
@@ -591,7 +591,7 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
CountersSize == CachedCntsSize) {
HostCountersBegin = CachedHostCnts;
CntsReused = 1;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Reusing cached counters section (%zu bytes)\n", CountersSize);
} else if (CountersSize > 0) {
HostCountersBegin = (char *)malloc(CountersSize);
@@ -601,7 +601,7 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
DataSize == CachedDataSize) {
HostDataBegin = CachedHostData;
DataReused = 1;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Reusing cached data section (%zu bytes)\n", DataSize);
} else if (DataSize > 0) {
HostDataBegin = (char *)malloc(DataSize);
@@ -611,7 +611,7 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
NamesSize == CachedNamesSize) {
HostNamesBegin = CachedHostNames;
NamesReused = 1;
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Reusing cached names section (%zu bytes)\n", NamesSize);
} else if (NamesSize > 0) {
HostNamesBegin = (char *)malloc(NamesSize);
@@ -660,7 +660,7 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
CachedNamesSize = NamesSize;
}
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Copied device sections: Counters=%zu, Data=%zu, Names=%zu\n",
CountersSize, DataSize, NamesSize);
@@ -742,14 +742,14 @@ static int ProcessDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
if (ret != 0) {
PROF_ERR("%s\n", "failed to write device profile using shared API");
- } else if (IsVerboseMode()) {
+ } else if (isVerboseMode()) {
PROF_NOTE("%s\n", "Successfully wrote device profile using shared API");
}
return ret;
}
-static int ProcessShadowVariable(void *ShadowVar, int TUIndex,
+static int processShadowVariable(void *ShadowVar, int TUIndex,
const char *Target) {
void *DevicePtrVar = nullptr;
if (hipGetSymbolAddress(&DevicePtrVar, ShadowVar) != 0) {
@@ -766,12 +766,12 @@ static int ProcessShadowVariable(void *ShadowVar, int TUIndex,
ShadowVar);
return -1;
}
- return ProcessDeviceOffloadPrf(DeviceOffloadPrf, TUIndex, Target);
+ return processDeviceOffloadPrf(DeviceOffloadPrf, TUIndex, Target);
}
/* Check if HIP runtime is available and loaded */
-static int IsHipAvailable(void) {
- EnsureHipLoaded();
+static int isHipAvailable(void) {
+ ensureHipLoaded();
return pHipMemcpy != nullptr && pHipGetSymbolAddress != nullptr;
}
@@ -783,7 +783,7 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
if (NumShadowVariables == 0 && NumDynamicModules == 0)
return 0;
- if (!IsHipAvailable())
+ if (!isHipAvailable())
return 0;
int Ret = 0;
@@ -796,16 +796,16 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
for (int Dev = 0; Dev < NumDevices; ++Dev) {
if (hipSetDevice(Dev) != 0) {
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Failed to set device %d, skipping\n", Dev);
continue;
}
const char *ArchName = getDeviceArchName(Dev);
- if (IsVerboseMode())
+ if (isVerboseMode())
PROF_NOTE("Collecting static profile data from device %d (%s)\n", Dev,
ArchName);
for (int i = 0; i < NumShadowVariables; ++i) {
- if (ProcessShadowVariable(OffloadShadowVariables[i], i, ArchName) != 0)
+ if (processShadowVariable(OffloadShadowVariables[i], i, ArchName) != 0)
Ret = -1;
}
}
>From c01b57218ac18eb5d456be48a2ae1628f28178d4 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Apr 2026 00:43:43 -0400
Subject: [PATCH 03/13] [PGO][AMDGPU] Apply follow-up review cleanups
Fold the remaining review cleanup into one change by replacing the sanitizer builtin workaround with compiler builtins, returning StringRef for CUID extraction, routing __llvm_profile_instrument_gpu through RuntimeLibcalls with an IR-level abstract-to-concrete lookup, deriving pointer types and sizes from the globals being described, and dropping redundant llvm:: qualifiers.
---
.../profile/InstrProfilingPlatformROCm.cpp | 47 ++++++---------
llvm/include/llvm/IR/RuntimeLibcalls.h | 22 +++++++
llvm/include/llvm/IR/RuntimeLibcalls.td | 11 +++-
llvm/lib/IR/RuntimeLibcalls.cpp | 12 ++++
.../Instrumentation/InstrProfiling.cpp | 58 +++++++++++++------
5 files changed, 101 insertions(+), 49 deletions(-)
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index 79646f17b29e2..040a44de2c5c3 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -12,17 +12,7 @@ extern "C" {
#include "InstrProfilingPort.h"
}
-// interception.h pulls in sanitizer_internal_defs.h, which normally includes
-// sanitizer_redefine_builtins.h. That uses inline asm to alias
-// memcpy/memmove/memset to __sanitizer_internal_* (see sanitizer_libc.cpp in
-// sanitizer_common). The instrumented *host* link for HIP (-fprofile-generate)
-// only pulls in libclang_rt.profile.a, not the full sanitizer_common objects
-// that define those symbols, so we get undefined references at link time. This
-// TU does not need the sanitizer builtin redirect; keep using libc
-// memcpy/memset.
-#define SANITIZER_COMMON_NO_REDEFINE_BUILTINS 1
#include "interception/interception.h"
-#undef SANITIZER_COMMON_NO_REDEFINE_BUILTINS
// 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>
@@ -127,7 +117,7 @@ static void ensureHipLoaded(void) {
Count = MAX_DEVICES;
HipDevicePropMinimal Prop;
for (int i = 0; i < Count; ++i) {
- memset(&Prop, 0, sizeof(Prop));
+ __builtin_memset(&Prop, 0, sizeof(Prop));
if (pHipGetDeviceProperties(&Prop, i) == 0) {
strncpy(DeviceArchNames[i], Prop.gcnArchName,
sizeof(DeviceArchNames[i]) - 1);
@@ -234,17 +224,18 @@ static const void *unwrapOffloadBundle(const void *Image) {
const char *Buf = (const char *)Image;
uint64_t NumEntries;
- memcpy(&NumEntries, Buf + sizeof(BundleMagic) - 1, sizeof(uint64_t));
+ __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;
- memcpy(&EntryOffset, Cursor, sizeof(EntryOffset));
+ __builtin_memcpy(&EntryOffset, Cursor, sizeof(EntryOffset));
Cursor += sizeof(EntryOffset);
- memcpy(&EntrySize, Cursor, sizeof(EntrySize));
+ __builtin_memcpy(&EntrySize, Cursor, sizeof(EntrySize));
Cursor += sizeof(EntrySize);
- memcpy(&IDSize, Cursor, sizeof(IDSize));
+ __builtin_memcpy(&IDSize, Cursor, sizeof(IDSize));
Cursor += sizeof(IDSize);
/* Skip the entry ID string. */
Cursor += IDSize;
@@ -691,16 +682,16 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
return -1;
}
char *ContiguousBuffer = ContiguousBuf.get();
- memset(ContiguousBuffer, 0, ContiguousBufferSize);
+ __builtin_memset(ContiguousBuffer, 0, ContiguousBufferSize);
char *BufDataBegin = ContiguousBuffer;
char *BufCountersBegin =
ContiguousBuffer + DataSize + PaddingBytesBeforeCounters;
char *BufNamesBegin = BufCountersBegin + CountersSize;
- memcpy(BufDataBegin, HostDataBegin, DataSize);
- memcpy(BufCountersBegin, HostCountersBegin, CountersSize);
- memcpy(BufNamesBegin, HostNamesBegin, NamesSize);
+ __builtin_memcpy(BufDataBegin, HostDataBegin, DataSize);
+ __builtin_memcpy(BufCountersBegin, HostCountersBegin, CountersSize);
+ __builtin_memcpy(BufNamesBegin, HostNamesBegin, NamesSize);
// Relocate CounterPtr in data records for file layout.
// CounterPtr is device-relative offset; adjust for file layout where
@@ -719,16 +710,16 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
ptrdiff_t NewRelativeOffset = DataSize + PaddingBytesBeforeCounters +
OffsetIntoCountersSection -
(i * sizeof(__llvm_profile_data));
- memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
- offsetof(__llvm_profile_data, CounterPtr),
- &NewRelativeOffset, sizeof(NewRelativeOffset));
+ __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, CounterPtr),
+ &NewRelativeOffset, sizeof(NewRelativeOffset));
}
- memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
- offsetof(__llvm_profile_data, BitmapPtr),
- 0,
- sizeof(RelocatedData[i].BitmapPtr) +
- sizeof(RelocatedData[i].FunctionPointer) +
- sizeof(RelocatedData[i].Values));
+ __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) +
+ offsetof(__llvm_profile_data, BitmapPtr),
+ 0,
+ sizeof(RelocatedData[i].BitmapPtr) +
+ sizeof(RelocatedData[i].FunctionPointer) +
+ sizeof(RelocatedData[i].Values));
}
char TUIndexStr[16];
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.h b/llvm/include/llvm/IR/RuntimeLibcalls.h
index fa092909d630d..3e517e0ae59d1 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.h
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.h
@@ -129,6 +129,21 @@ struct RuntimeLibcallsInfo {
AvailableLibcallImpls.set(Impl);
}
+ /// Return the first available concrete impl that provides the abstract
+ /// libcall \p LC for the current module's target, or RTLIB::Unsupported
+ /// if no available impl provides it. Mirrors codegen-side
+ /// LibcallLoweringInfo::getLibcallImpl, but available in IR passes that
+ /// don't have a TargetSubtargetInfo.
+ RTLIB::LibcallImpl getLibcallImpl(RTLIB::Libcall LC) const {
+ return LibcallImpls[LC];
+ }
+
+ /// Convenience: return the entry-point name for the abstract libcall \p LC,
+ /// or an empty StringRef if no impl is available.
+ StringRef getLibcallName(RTLIB::Libcall LC) const {
+ return getLibcallImplName(getLibcallImpl(LC));
+ }
+
/// Check if a function name is a recognized runtime call of any kind. This
/// does not consider if this call is available for any current compilation,
/// just that it is a known call somewhere. This returns the set of all
@@ -178,6 +193,13 @@ struct RuntimeLibcallsInfo {
/// implementation.;
CallingConv::ID LibcallImplCallingConvs[RTLIB::NumLibcallImpls] = {};
+ /// Cache mapping each abstract Libcall to the first available concrete impl
+ /// for the current target. Populated at the end of the constructor (after
+ /// initLibcalls and any vector-library setAvailable calls have run). First
+ /// available impl wins, matching codegen-side LibcallLoweringInfo.
+ RTLIB::LibcallImpl LibcallImpls[RTLIB::UNKNOWN_LIBCALL + 1] = {
+ RTLIB::Unsupported};
+
/// Names of concrete implementations of runtime calls. e.g. __ashlsi3 for
/// SHL_I32
LLVM_ABI static const char RuntimeLibcallImplNameTableStorage[];
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.td b/llvm/include/llvm/IR/RuntimeLibcalls.td
index cbdc48a9a717f..10df06a600fa0 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.td
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.td
@@ -504,6 +504,9 @@ def RETURN_ADDRESS : RuntimeLibcall;
def CLEAR_CACHE : RuntimeLibcall;
def RISCV_FLUSH_ICACHE : RuntimeLibcall;
+// Profile instrumentation runtime
+def INSTR_PROF_INSTRUMENT_GPU : RuntimeLibcall;
+
// Mips16 calls
def MIPS16_RET_DC : RuntimeLibcall;
def MIPS16_RET_DF : RuntimeLibcall;
@@ -1157,6 +1160,10 @@ def __llvm_deoptimize : RuntimeLibcallImpl<DEOPTIMIZE>;
// Clear cache
def __clear_cache : RuntimeLibcallImpl<CLEAR_CACHE>;
+// Profile instrumentation runtime
+def __llvm_profile_instrument_gpu :
+ RuntimeLibcallImpl<INSTR_PROF_INSTRUMENT_GPU>;
+
//--------------------------------------------------------------------
// libm
//--------------------------------------------------------------------
@@ -2241,8 +2248,8 @@ def WindowsARM64ECSystemLibrary
def isAMDGPU : RuntimeLibcallPredicate<"TT.isAMDGPU()">;
-// No calls.
-def AMDGPUSystemLibrary : SystemRuntimeLibrary<isAMDGPU, (add)>;
+def AMDGPUSystemLibrary :
+ SystemRuntimeLibrary<isAMDGPU, (add __llvm_profile_instrument_gpu)>;
//===----------------------------------------------------------------------===//
// ARM Runtime Libcalls
diff --git a/llvm/lib/IR/RuntimeLibcalls.cpp b/llvm/lib/IR/RuntimeLibcalls.cpp
index d72277fa2b179..a9d4ce8f478e6 100644
--- a/llvm/lib/IR/RuntimeLibcalls.cpp
+++ b/llvm/lib/IR/RuntimeLibcalls.cpp
@@ -99,6 +99,18 @@ RuntimeLibcallsInfo::RuntimeLibcallsInfo(const Triple &TT,
default:
break;
}
+
+ // Populate the abstract Libcall -> first available LibcallImpl cache so that
+ // IR passes (which don't have a TargetSubtargetInfo) can resolve a libcall
+ // name from its abstract enum the same way codegen does. Matches the
+ // "first available impl wins" policy used by LibcallLoweringInfo.
+ for (RTLIB::LibcallImpl Impl : RTLIB::libcall_impls()) {
+ if (!isAvailable(Impl))
+ continue;
+ RTLIB::Libcall LC = getLibcallFromImpl(Impl);
+ if (LibcallImpls[LC] == RTLIB::Unsupported)
+ LibcallImpls[LC] = Impl;
+ }
}
RuntimeLibcallsInfo::RuntimeLibcallsInfo(const Module &M)
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index f930b649df1b1..5d4317cf18723 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -22,6 +22,7 @@
#include "llvm/Analysis/BranchProbabilityInfo.h"
#include "llvm/Analysis/CFG.h"
#include "llvm/Analysis/LoopInfo.h"
+#include "llvm/Analysis/RuntimeLibcallInfo.h"
#include "llvm/Analysis/TargetLibraryInfo.h"
#include "llvm/Frontend/Offloading/Utility.h"
#include "llvm/IR/Attributes.h"
@@ -45,6 +46,7 @@
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
+#include "llvm/IR/RuntimeLibcalls.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
@@ -247,14 +249,16 @@ static bool profDataReferencedByCode(const Module &M) {
// Extract CUID (Compilation Unit ID) from the module.
// HIP/CUDA modules have a global variable __hip_cuid_<hash> that uniquely
-// identifies each translation unit. Returns empty string if not found.
-static std::string getCUIDFromModule(const Module &M) {
+// identifies each translation unit. Returns an empty StringRef if not found.
+// The returned StringRef points into the GlobalVariable's name, which is
+// owned by the Module and stable for the Module's lifetime.
+static StringRef getCUIDFromModule(const Module &M) {
for (const GlobalVariable &GV : M.globals()) {
if (!GV.hasExternalLinkage())
continue;
StringRef Name = GV.getName();
if (Name.consume_front("__hip_cuid_"))
- return Name.str();
+ return Name;
}
return "";
}
@@ -263,9 +267,10 @@ class InstrLowerer final {
public:
InstrLowerer(Module &M, const InstrProfOptions &Options,
std::function<const TargetLibraryInfo &(Function &F)> GetTLI,
- bool IsCS)
+ const RTLIB::RuntimeLibcallsInfo &RTLCI, bool IsCS)
: M(M), Options(Options), TT(M.getTargetTriple()), IsCS(IsCS),
- GetTLI(GetTLI), DataReferencedByCode(profDataReferencedByCode(M)) {}
+ GetTLI(GetTLI), RTLCI(RTLCI),
+ DataReferencedByCode(profDataReferencedByCode(M)) {}
bool lower();
@@ -278,6 +283,9 @@ class InstrLowerer final {
std::function<const TargetLibraryInfo &(Function &F)> GetTLI;
+ // Runtime libcall registry (e.g. __llvm_profile_instrument_gpu).
+ const RTLIB::RuntimeLibcallsInfo &RTLCI;
+
const bool DataReferencedByCode;
struct PerFunctionProfileData {
@@ -714,7 +722,9 @@ PreservedAnalyses InstrProfilingLoweringPass::run(Module &M,
auto GetTLI = [&FAM](Function &F) -> TargetLibraryInfo & {
return FAM.getResult<TargetLibraryAnalysis>(F);
};
- InstrLowerer Lowerer(M, Options, GetTLI, IsCS);
+ const RTLIB::RuntimeLibcallsInfo &RTLCI =
+ AM.getResult<RuntimeLibraryAnalysis>(M);
+ InstrLowerer Lowerer(M, Options, GetTLI, RTLCI, IsCS);
if (!Lowerer.lower())
return PreservedAnalyses::all();
@@ -1301,8 +1311,18 @@ void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
{PtrTy, PtrTy, Int64Ty}, false);
- FunctionCallee IncrFn =
- M.getOrInsertFunction("__llvm_profile_instrument_gpu", CalleeTy);
+ // Look up the runtime entry-point name through the RuntimeLibcalls
+ // registry instead of hardcoding the string. The abstract libcall is
+ // registered in llvm/include/llvm/IR/RuntimeLibcalls.td and a concrete
+ // impl is added to AMDGPUSystemLibrary; the consumer side is in
+ // compiler-rt's clang_rt.profile.
+ RTLIB::LibcallImpl IncrImpl =
+ RTLCI.getLibcallImpl(RTLIB::INSTR_PROF_INSTRUMENT_GPU);
+ assert(IncrImpl != RTLIB::Unsupported &&
+ "RuntimeLibcalls.td must register INSTR_PROF_INSTRUMENT_GPU "
+ "for any target that reaches lowerIncrementAMDGPU");
+ StringRef IncrFnName = RTLCI.getLibcallImplName(IncrImpl);
+ FunctionCallee IncrFn = M.getOrInsertFunction(IncrFnName, CalleeTy);
Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
Inc->eraseFromParent();
@@ -2173,9 +2193,10 @@ void InstrLowerer::emitNameData() {
ConstantDataArray::getString(Ctx, StringRef(CompressedNameStr), false);
std::string NamesVarName = std::string(getInstrProfNamesVarName());
if (isGPUProfTarget(M)) {
- std::string CUID = CachedCUID.empty() ? getCUIDFromModule(M) : CachedCUID;
+ StringRef CUID =
+ CachedCUID.empty() ? getCUIDFromModule(M) : StringRef(CachedCUID);
if (!CUID.empty())
- NamesVarName = NamesVarName + "_" + CUID;
+ NamesVarName = (Twine(NamesVarName) + "_" + CUID).str();
}
NamesVar =
new GlobalVariable(M, NamesVal->getType(), true,
@@ -2455,7 +2476,7 @@ void InstrLowerer::createProfileSectionSymbols() {
// The host reads through this indirection: hipGetSymbolAddress gives the
// pointer global's device address, then one DtoH copy yields the sections
// struct address, then another DtoH copy reads the actual sections.
- auto *PtrTy = PointerType::get(Ctx, AS);
+ auto *PtrTy = SectionsGV->getType();
auto *PtrInit =
ConstantExpr::getPointerBitCastOrAddrSpaceCast(SectionsGV, PtrTy);
std::string PtrName = "__llvm_offload_prf_" + CachedCUID;
@@ -2470,7 +2491,8 @@ void InstrLowerer::createHIPDeviceVariableRegistration() {
if (isGPUProfTarget(M))
return;
- std::string CUID = CachedCUID.empty() ? getCUIDFromModule(M) : CachedCUID;
+ StringRef CUID =
+ CachedCUID.empty() ? getCUIDFromModule(M) : StringRef(CachedCUID);
if (CUID.empty())
return;
@@ -2478,7 +2500,7 @@ void InstrLowerer::createHIPDeviceVariableRegistration() {
auto *VoidTy = Type::getVoidTy(Ctx);
auto *VoidPtrTy = PointerType::getUnqual(Ctx);
- std::string OffloadPrfName = "__llvm_offload_prf_" + CUID;
+ std::string OffloadPrfName = ("__llvm_offload_prf_" + CUID).str();
auto *OffloadPrfShadow = new GlobalVariable(
M, VoidPtrTy, /*isConstant=*/false, GlobalValue::ExternalLinkage,
ConstantPointerNull::get(cast<PointerType>(VoidPtrTy)), OffloadPrfName);
@@ -2494,7 +2516,7 @@ void InstrLowerer::createHIPDeviceVariableRegistration() {
// linker wrapper generates __hipRegisterVar in the final module ctor.
llvm::offloading::emitOffloadingEntry(
M, llvm::object::OffloadKind::OFK_HIP, OffloadPrfShadow, OffloadPrfName,
- M.getDataLayout().getPointerSize(),
+ M.getDataLayout().getPointerSize(VoidPtrTy->getPointerAddressSpace()),
llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
auto *CtorFn = Function::Create(FunctionType::get(VoidTy, false),
@@ -2535,8 +2557,7 @@ void InstrLowerer::createHIPDeviceVariableRegistration() {
}
if (!FatbinHandleGV) {
- LLVM_DEBUG(llvm::dbgs()
- << "store of __hipRegisterFatBinary call not found\n");
+ LLVM_DEBUG(dbgs() << "store of __hipRegisterFatBinary call not found\n");
}
// Insert the new registration just before the ctor’s return
@@ -2548,9 +2569,8 @@ void InstrLowerer::createHIPDeviceVariableRegistration() {
return;
IRBuilder<> Builder(RetInst);
- LLVM_DEBUG(
- llvm::dbgs() << "Found __hip_module_ctor, registering anchors for CUID="
- << CUID << "\n");
+ LLVM_DEBUG(dbgs() << "Found __hip_module_ctor, registering anchors for CUID="
+ << CUID << "\n");
auto *Int32Ty = Type::getInt32Ty(Ctx);
auto *Int64Ty = Type::getInt64Ty(Ctx);
>From 73cfd15f1f2f8fb49c6a8ffd1be91e44f57e3962 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Apr 2026 12:51:47 -0400
Subject: [PATCH 04/13] [PGO][AMDGPU] Drop the IR RuntimeLibcalls helper
Remove the IR-side RuntimeLibcallsInfo abstract-to-concrete lookup helper and switch InstrProfiling back to the concrete __llvm_profile_instrument_gpu impl while keeping the RuntimeLibcalls.td entry that centralizes the name for AMDGPU. This avoids adding new lower-level runtime-libcall infrastructure for an IR-pass-only use case and matches the existing pattern where subtarget-aware libcall selection lives in codegen passes, not llvm/lib/Transforms.
---
llvm/include/llvm/IR/RuntimeLibcalls.h | 22 -------------------
llvm/lib/IR/RuntimeLibcalls.cpp | 12 ----------
.../Instrumentation/InstrProfiling.cpp | 17 +++++++-------
3 files changed, 8 insertions(+), 43 deletions(-)
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.h b/llvm/include/llvm/IR/RuntimeLibcalls.h
index 3e517e0ae59d1..fa092909d630d 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.h
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.h
@@ -129,21 +129,6 @@ struct RuntimeLibcallsInfo {
AvailableLibcallImpls.set(Impl);
}
- /// Return the first available concrete impl that provides the abstract
- /// libcall \p LC for the current module's target, or RTLIB::Unsupported
- /// if no available impl provides it. Mirrors codegen-side
- /// LibcallLoweringInfo::getLibcallImpl, but available in IR passes that
- /// don't have a TargetSubtargetInfo.
- RTLIB::LibcallImpl getLibcallImpl(RTLIB::Libcall LC) const {
- return LibcallImpls[LC];
- }
-
- /// Convenience: return the entry-point name for the abstract libcall \p LC,
- /// or an empty StringRef if no impl is available.
- StringRef getLibcallName(RTLIB::Libcall LC) const {
- return getLibcallImplName(getLibcallImpl(LC));
- }
-
/// Check if a function name is a recognized runtime call of any kind. This
/// does not consider if this call is available for any current compilation,
/// just that it is a known call somewhere. This returns the set of all
@@ -193,13 +178,6 @@ struct RuntimeLibcallsInfo {
/// implementation.;
CallingConv::ID LibcallImplCallingConvs[RTLIB::NumLibcallImpls] = {};
- /// Cache mapping each abstract Libcall to the first available concrete impl
- /// for the current target. Populated at the end of the constructor (after
- /// initLibcalls and any vector-library setAvailable calls have run). First
- /// available impl wins, matching codegen-side LibcallLoweringInfo.
- RTLIB::LibcallImpl LibcallImpls[RTLIB::UNKNOWN_LIBCALL + 1] = {
- RTLIB::Unsupported};
-
/// Names of concrete implementations of runtime calls. e.g. __ashlsi3 for
/// SHL_I32
LLVM_ABI static const char RuntimeLibcallImplNameTableStorage[];
diff --git a/llvm/lib/IR/RuntimeLibcalls.cpp b/llvm/lib/IR/RuntimeLibcalls.cpp
index a9d4ce8f478e6..d72277fa2b179 100644
--- a/llvm/lib/IR/RuntimeLibcalls.cpp
+++ b/llvm/lib/IR/RuntimeLibcalls.cpp
@@ -99,18 +99,6 @@ RuntimeLibcallsInfo::RuntimeLibcallsInfo(const Triple &TT,
default:
break;
}
-
- // Populate the abstract Libcall -> first available LibcallImpl cache so that
- // IR passes (which don't have a TargetSubtargetInfo) can resolve a libcall
- // name from its abstract enum the same way codegen does. Matches the
- // "first available impl wins" policy used by LibcallLoweringInfo.
- for (RTLIB::LibcallImpl Impl : RTLIB::libcall_impls()) {
- if (!isAvailable(Impl))
- continue;
- RTLIB::Libcall LC = getLibcallFromImpl(Impl);
- if (LibcallImpls[LC] == RTLIB::Unsupported)
- LibcallImpls[LC] = Impl;
- }
}
RuntimeLibcallsInfo::RuntimeLibcallsInfo(const Module &M)
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index 5d4317cf18723..744a842507cc1 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -1312,16 +1312,15 @@ void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
{PtrTy, PtrTy, Int64Ty}, false);
// Look up the runtime entry-point name through the RuntimeLibcalls
- // registry instead of hardcoding the string. The abstract libcall is
- // registered in llvm/include/llvm/IR/RuntimeLibcalls.td and a concrete
- // impl is added to AMDGPUSystemLibrary; the consumer side is in
+ // registry instead of hardcoding the string. The concrete impl is
+ // registered for AMDGPU in llvm/include/llvm/IR/RuntimeLibcalls.td; the
+ // consumer side is in
// compiler-rt's clang_rt.profile.
- RTLIB::LibcallImpl IncrImpl =
- RTLCI.getLibcallImpl(RTLIB::INSTR_PROF_INSTRUMENT_GPU);
- assert(IncrImpl != RTLIB::Unsupported &&
- "RuntimeLibcalls.td must register INSTR_PROF_INSTRUMENT_GPU "
- "for any target that reaches lowerIncrementAMDGPU");
- StringRef IncrFnName = RTLCI.getLibcallImplName(IncrImpl);
+ assert(RTLCI.isAvailable(RTLIB::impl___llvm_profile_instrument_gpu) &&
+ "RuntimeLibcalls.td must register __llvm_profile_instrument_gpu "
+ "for AMDGPU");
+ StringRef IncrFnName =
+ RTLCI.getLibcallImplName(RTLIB::impl___llvm_profile_instrument_gpu);
FunctionCallee IncrFn = M.getOrInsertFunction(IncrFnName, CalleeTy);
Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
>From 8126bb1ab30903c106406af273c927b44f3183e6 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Apr 2026 14:02:32 -0400
Subject: [PATCH 05/13] [PGO][AMDGPU] Use the existing InstrProf callback
naming pattern
Drop the RuntimeLibcalls entry for __llvm_profile_instrument_gpu and reach
the runtime entry-point through the existing profile naming helper. The
INSTR_PROF_INSTRUMENT_GPU_FUNC macro added in #196538 is exposed via a small
getInstrProfInstrumentGPUFuncName() accessor in InstrProf.h, which the
AMDGPU lowering in InstrProfiling.cpp now uses. This avoids mixing a
codegen-style runtime-libcall convention into an IR-pass-only call site
while keeping the symbol name centralized in the profile runtime naming
layer.
---
llvm/include/llvm/IR/RuntimeLibcalls.td | 11 ++------
llvm/include/llvm/ProfileData/InstrProf.h | 5 ++++
.../Instrumentation/InstrProfiling.cpp | 27 ++++---------------
.../Instrumentation/PGOInstrumentation.cpp | 3 ++-
4 files changed, 14 insertions(+), 32 deletions(-)
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.td b/llvm/include/llvm/IR/RuntimeLibcalls.td
index 10df06a600fa0..cbdc48a9a717f 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.td
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.td
@@ -504,9 +504,6 @@ def RETURN_ADDRESS : RuntimeLibcall;
def CLEAR_CACHE : RuntimeLibcall;
def RISCV_FLUSH_ICACHE : RuntimeLibcall;
-// Profile instrumentation runtime
-def INSTR_PROF_INSTRUMENT_GPU : RuntimeLibcall;
-
// Mips16 calls
def MIPS16_RET_DC : RuntimeLibcall;
def MIPS16_RET_DF : RuntimeLibcall;
@@ -1160,10 +1157,6 @@ def __llvm_deoptimize : RuntimeLibcallImpl<DEOPTIMIZE>;
// Clear cache
def __clear_cache : RuntimeLibcallImpl<CLEAR_CACHE>;
-// Profile instrumentation runtime
-def __llvm_profile_instrument_gpu :
- RuntimeLibcallImpl<INSTR_PROF_INSTRUMENT_GPU>;
-
//--------------------------------------------------------------------
// libm
//--------------------------------------------------------------------
@@ -2248,8 +2241,8 @@ def WindowsARM64ECSystemLibrary
def isAMDGPU : RuntimeLibcallPredicate<"TT.isAMDGPU()">;
-def AMDGPUSystemLibrary :
- SystemRuntimeLibrary<isAMDGPU, (add __llvm_profile_instrument_gpu)>;
+// No calls.
+def AMDGPUSystemLibrary : SystemRuntimeLibrary<isAMDGPU, (add)>;
//===----------------------------------------------------------------------===//
// ARM Runtime Libcalls
diff --git a/llvm/include/llvm/ProfileData/InstrProf.h b/llvm/include/llvm/ProfileData/InstrProf.h
index dffc58281c2d9..b7f917ec39b70 100644
--- a/llvm/include/llvm/ProfileData/InstrProf.h
+++ b/llvm/include/llvm/ProfileData/InstrProf.h
@@ -122,6 +122,11 @@ inline StringRef getInstrProfValueProfMemOpFuncName() {
/// Return the prefix of the name of the variables to function as a filter.
inline StringRef getInstrProfVarPrefix() { return "__prof"; }
+/// Return the name of the GPU wave-cooperative counter increment helper.
+inline StringRef getInstrProfInstrumentGPUFuncName() {
+ return INSTR_PROF_INSTRUMENT_GPU_FUNC_STR;
+}
+
/// Return the name prefix of variables containing instrumented function names.
inline StringRef getInstrProfNameVarPrefix() { return "__profn_"; }
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index 744a842507cc1..2ed6b2f917b12 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -22,7 +22,6 @@
#include "llvm/Analysis/BranchProbabilityInfo.h"
#include "llvm/Analysis/CFG.h"
#include "llvm/Analysis/LoopInfo.h"
-#include "llvm/Analysis/RuntimeLibcallInfo.h"
#include "llvm/Analysis/TargetLibraryInfo.h"
#include "llvm/Frontend/Offloading/Utility.h"
#include "llvm/IR/Attributes.h"
@@ -46,7 +45,6 @@
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
-#include "llvm/IR/RuntimeLibcalls.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
@@ -267,10 +265,9 @@ class InstrLowerer final {
public:
InstrLowerer(Module &M, const InstrProfOptions &Options,
std::function<const TargetLibraryInfo &(Function &F)> GetTLI,
- const RTLIB::RuntimeLibcallsInfo &RTLCI, bool IsCS)
+ bool IsCS)
: M(M), Options(Options), TT(M.getTargetTriple()), IsCS(IsCS),
- GetTLI(GetTLI), RTLCI(RTLCI),
- DataReferencedByCode(profDataReferencedByCode(M)) {}
+ GetTLI(GetTLI), DataReferencedByCode(profDataReferencedByCode(M)) {}
bool lower();
@@ -283,9 +280,6 @@ class InstrLowerer final {
std::function<const TargetLibraryInfo &(Function &F)> GetTLI;
- // Runtime libcall registry (e.g. __llvm_profile_instrument_gpu).
- const RTLIB::RuntimeLibcallsInfo &RTLCI;
-
const bool DataReferencedByCode;
struct PerFunctionProfileData {
@@ -722,9 +716,7 @@ PreservedAnalyses InstrProfilingLoweringPass::run(Module &M,
auto GetTLI = [&FAM](Function &F) -> TargetLibraryInfo & {
return FAM.getResult<TargetLibraryAnalysis>(F);
};
- const RTLIB::RuntimeLibcallsInfo &RTLCI =
- AM.getResult<RuntimeLibraryAnalysis>(M);
- InstrLowerer Lowerer(M, Options, GetTLI, RTLCI, IsCS);
+ InstrLowerer Lowerer(M, Options, GetTLI, IsCS);
if (!Lowerer.lower())
return PreservedAnalyses::all();
@@ -1311,17 +1303,8 @@ void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
{PtrTy, PtrTy, Int64Ty}, false);
- // Look up the runtime entry-point name through the RuntimeLibcalls
- // registry instead of hardcoding the string. The concrete impl is
- // registered for AMDGPU in llvm/include/llvm/IR/RuntimeLibcalls.td; the
- // consumer side is in
- // compiler-rt's clang_rt.profile.
- assert(RTLCI.isAvailable(RTLIB::impl___llvm_profile_instrument_gpu) &&
- "RuntimeLibcalls.td must register __llvm_profile_instrument_gpu "
- "for AMDGPU");
- StringRef IncrFnName =
- RTLCI.getLibcallImplName(RTLIB::impl___llvm_profile_instrument_gpu);
- FunctionCallee IncrFn = M.getOrInsertFunction(IncrFnName, CalleeTy);
+ FunctionCallee IncrFn =
+ M.getOrInsertFunction(getInstrProfInstrumentGPUFuncName(), CalleeTy);
Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
Inc->eraseFromParent();
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index 404c47d5183f1..88677c641646e 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -2408,7 +2408,8 @@ void llvm::setProfMetadata(Instruction *TI, ArrayRef<uint64_t> EdgeCounts,
uint64_t MaxCount) {
auto Weights = downscaleWeights(EdgeCounts, MaxCount);
- LLVM_DEBUG(dbgs() << "Weight is: "; for (const auto &W : Weights) {
+ LLVM_DEBUG(dbgs() << "Weight is: "; for (const auto &W
+ : Weights) {
dbgs() << W << " ";
} dbgs() << "\n";);
>From e48e017d21830fff75f2b508ae3f431e67200599 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 8 May 2026 15:41:47 -0400
Subject: [PATCH 06/13] [PGO][AMDGPU] Drop unrelated drive-by changes
Three small cleanups split out from the rest of the series so the per-file
delta is only what the AMDGPU PGO feature actually needs:
- PGOInstrumentation.cpp: revert two pre-existing `Triple TT(M.getTargetTriple());`
copies to themselves; the new use inside isValueProfilingDisabled() is kept.
- InstrProfiling.cpp: collapse the unused `GlobalValue *DataVar = Data; Constant
*DataAddr = Data;` aliases (and the matching field/parameter widening on
PerFunctionProfileData::DataVar and lowerValueProfileInst) back to using
`Data` / `GlobalVariable *` directly. The widening was leftover from an
earlier iteration and was reversed at every use via cast<GlobalVariable>.
- InstrProfiling.cpp: restore the WHY comment on the NVPTX absolute-pointer
branch in createDataVariable; drop a redundant comment on getCounterAddress
that just narrated the AMDGPU early-exit on the next function.
---
.../Instrumentation/InstrProfiling.cpp | 39 +++++++++----------
.../Instrumentation/PGOInstrumentation.cpp | 4 +-
2 files changed, 20 insertions(+), 23 deletions(-)
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index 2ed6b2f917b12..aba5b1331e2a8 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -285,7 +285,7 @@ class InstrLowerer final {
struct PerFunctionProfileData {
uint32_t NumValueSites[IPVK_Last + 1] = {};
GlobalVariable *RegionCounters = nullptr;
- GlobalValue *DataVar = nullptr;
+ GlobalVariable *DataVar = nullptr;
GlobalVariable *RegionBitmaps = nullptr;
uint32_t NumBitmapBytes = 0;
@@ -1102,7 +1102,7 @@ void InstrLowerer::lowerValueProfileInst(InstrProfValueProfileInst *Ind) {
assert(It != ProfileDataMap.end() && It->second.DataVar &&
"value profiling detected in function with no counter increment");
- GlobalValue *DataVar = It->second.DataVar;
+ GlobalVariable *DataVar = It->second.DataVar;
uint64_t ValueKind = Ind->getValueKind()->getZExtValue();
uint64_t Index = Ind->getIndex()->getZExtValue();
for (uint32_t Kind = IPVK_First; Kind < ValueKind; ++Kind)
@@ -1114,7 +1114,7 @@ void InstrLowerer::lowerValueProfileInst(InstrProfValueProfileInst *Ind) {
CallInst *Call = nullptr;
auto *TLI = &GetTLI(*Ind->getFunction());
auto *NormalizedDataVarPtr = ConstantExpr::getPointerBitCastOrAddrSpaceCast(
- cast<Constant>(DataVar), PointerType::get(M.getContext(), 0));
+ DataVar, PointerType::get(M.getContext(), 0));
// To support value profiling calls within Windows exception handlers, funclet
// information contained within operand bundles needs to be copied over to
@@ -1164,8 +1164,6 @@ GlobalVariable *InstrLowerer::getOrCreateBiasVar(StringRef VarName) {
}
Value *InstrLowerer::getCounterAddress(InstrProfCntrInstBase *I) {
- // Note: For AMDGPU targets, lowerIncrementAMDGPU handles counter addressing
- // directly. This function is called for non-AMDGPU targets.
auto *Counters = getOrCreateRegionCounters(I);
IRBuilder<> Builder(I);
@@ -1952,8 +1950,6 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
Visibility = GlobalValue::ProtectedVisibility;
auto *Data =
new GlobalVariable(M, DataTy, false, Linkage, nullptr, DataVarName);
- GlobalValue *DataVar = Data;
- Constant *DataAddr = Data;
Constant *RelativeCounterPtr;
GlobalVariable *BitmapPtr = PD.RegionBitmaps;
@@ -1967,6 +1963,9 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
if (BitmapPtr != nullptr)
RelativeBitmapPtr = ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy);
} else if (TT.isNVPTX()) {
+ // The NVPTX target cannot handle self-referencing constant expressions in
+ // global initializers at all. Use absolute pointers and have the runtime
+ // registration convert them to relative offsets.
DataSectionKind = IPSK_data;
RelativeCounterPtr = ConstantExpr::getPtrToInt(CounterPtr, IntPtrTy);
} else {
@@ -1975,36 +1974,34 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
DataSectionKind = IPSK_data;
RelativeCounterPtr =
ConstantExpr::getSub(ConstantExpr::getPtrToInt(CounterPtr, IntPtrTy),
- ConstantExpr::getPtrToInt(DataAddr, IntPtrTy));
+ ConstantExpr::getPtrToInt(Data, IntPtrTy));
if (BitmapPtr != nullptr)
RelativeBitmapPtr =
ConstantExpr::getSub(ConstantExpr::getPtrToInt(BitmapPtr, IntPtrTy),
- ConstantExpr::getPtrToInt(DataAddr, IntPtrTy));
+ ConstantExpr::getPtrToInt(Data, IntPtrTy));
}
Constant *DataVals[] = {
#define INSTR_PROF_DATA(Type, LLVMType, Name, Init) Init,
#include "llvm/ProfileData/InstrProfData.inc"
};
- auto *DataInit = ConstantStruct::get(DataTy, DataVals);
+ Data->setInitializer(ConstantStruct::get(DataTy, DataVals));
- auto *DataGV = cast<GlobalVariable>(DataVar);
- DataGV->setInitializer(DataInit);
- DataGV->setVisibility(Visibility);
- DataGV->setSection(
+ Data->setVisibility(Visibility);
+ Data->setSection(
getInstrProfSectionName(DataSectionKind, TT.getObjectFormat()));
- DataGV->setAlignment(Align(INSTR_PROF_DATA_ALIGNMENT));
- if (isGPUProfTarget(M) && !DataGV->hasComdat()) {
- DataGV->setComdat(M.getOrInsertComdat(CntsVarName));
- DataGV->setLinkage(GlobalValue::LinkOnceODRLinkage);
+ Data->setAlignment(Align(INSTR_PROF_DATA_ALIGNMENT));
+ if (isGPUProfTarget(M) && !Data->hasComdat()) {
+ Data->setComdat(M.getOrInsertComdat(CntsVarName));
+ Data->setLinkage(GlobalValue::LinkOnceODRLinkage);
} else {
- maybeSetComdat(DataGV, Fn, CntsVarName);
+ maybeSetComdat(Data, Fn, CntsVarName);
}
- PD.DataVar = DataVar;
+ PD.DataVar = Data;
// Mark the data variable as used so that it isn't stripped out.
- CompilerUsedVars.push_back(DataVar);
+ CompilerUsedVars.push_back(Data);
// Now that the linkage set by the FE has been passed to the data and counter
// variables, reset Name variable's linkage and visibility to private so that
// it can be removed later by the compiler.
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index 88677c641646e..e0bfccd7cd54d 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -472,7 +472,7 @@ createIRLevelProfileFlagVar(Module &M,
Constant::getIntegerValue(IntTy64, APInt(64, ProfileVersion)), VarName);
IRLevelVersionVariable->setVisibility(GlobalValue::HiddenVisibility);
- const Triple &TT = M.getTargetTriple();
+ Triple TT(M.getTargetTriple());
if (TT.supportsCOMDAT()) {
IRLevelVersionVariable->setLinkage(GlobalValue::ExternalLinkage);
IRLevelVersionVariable->setComdat(M.getOrInsertComdat(VarName));
@@ -1961,7 +1961,7 @@ static bool InstrumentAllFunctions(
if (InstrumentationType == PGOInstrumentationType::FDO)
createIRLevelProfileFlagVar(M, InstrumentationType);
- const Triple &TT = M.getTargetTriple();
+ Triple TT(M.getTargetTriple());
LLVMContext &Ctx = M.getContext();
if (!TT.isOSBinFormatELF() && EnableVTableValueProfiling)
Ctx.diagnose(DiagnosticInfoPGOProfile(
>From b855f1f334e1cd1373bde443d7962fa574a844ec Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Sun, 17 May 2026 10:45:57 -0400
Subject: [PATCH 07/13] [PGO][AMDGPU] Look up __llvm_profile_instrument_gpu via
RuntimeLibcalls
Route lowerIncrementAMDGPU through the static
RuntimeLibcallsInfo::getLibcallImplName accessor, and register
PROFILE_INSTRUMENT_GPU plus the AMDGPU impl in RuntimeLibcalls.td.
---
llvm/include/llvm/IR/RuntimeLibcalls.td | 9 +++++++--
llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp | 5 ++++-
2 files changed, 11 insertions(+), 3 deletions(-)
diff --git a/llvm/include/llvm/IR/RuntimeLibcalls.td b/llvm/include/llvm/IR/RuntimeLibcalls.td
index cbdc48a9a717f..37bad559f49e7 100644
--- a/llvm/include/llvm/IR/RuntimeLibcalls.td
+++ b/llvm/include/llvm/IR/RuntimeLibcalls.td
@@ -500,6 +500,9 @@ def DEOPTIMIZE : RuntimeLibcall;
// Return address
def RETURN_ADDRESS : RuntimeLibcall;
+// GPU profiling
+def PROFILE_INSTRUMENT_GPU : RuntimeLibcall;
+
// Clear cache
def CLEAR_CACHE : RuntimeLibcall;
def RISCV_FLUSH_ICACHE : RuntimeLibcall;
@@ -2241,8 +2244,10 @@ def WindowsARM64ECSystemLibrary
def isAMDGPU : RuntimeLibcallPredicate<"TT.isAMDGPU()">;
-// No calls.
-def AMDGPUSystemLibrary : SystemRuntimeLibrary<isAMDGPU, (add)>;
+def __llvm_profile_instrument_gpu : RuntimeLibcallImpl<PROFILE_INSTRUMENT_GPU>;
+
+def AMDGPUSystemLibrary
+ : SystemRuntimeLibrary<isAMDGPU, (add __llvm_profile_instrument_gpu)>;
//===----------------------------------------------------------------------===//
// ARM Runtime Libcalls
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index aba5b1331e2a8..5bb205ed3c46a 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -45,6 +45,7 @@
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
+#include "llvm/IR/RuntimeLibcalls.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
@@ -1302,7 +1303,9 @@ void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
{PtrTy, PtrTy, Int64Ty}, false);
FunctionCallee IncrFn =
- M.getOrInsertFunction(getInstrProfInstrumentGPUFuncName(), CalleeTy);
+ M.getOrInsertFunction(RTLIB::RuntimeLibcallsInfo::getLibcallImplName(
+ RTLIB::impl___llvm_profile_instrument_gpu),
+ CalleeTy);
Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
Inc->eraseFromParent();
>From 94eb4e4857cd294080c788e3d299f411b09ad327 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Sun, 17 May 2026 14:41:32 -0400
Subject: [PATCH 08/13] [compiler-rt][profile] Make HIP runtime loading and
dynamic module list thread-safe
The first-time HIP runtime symbol resolver and the dynamic module
registration array were racy: two threads could enter ensureHipLoaded at
once, and concurrent registrations could realloc the DynamicModules
buffer underneath a thread that was reading it. Wrap the resolver in
pthread_once (POSIX) or INIT_ONCE (Windows), and guard the register,
unregister, and collect paths that touch DynamicModules with a pthread
mutex (POSIX) or CRITICAL_SECTION (Windows).
Addresses jhuber6 review on PR #177665.
---
.../profile/InstrProfilingPlatformROCm.cpp | 81 ++++++++++++++++---
1 file changed, 71 insertions(+), 10 deletions(-)
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index 040a44de2c5c3..24f10cd02e996 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -20,6 +20,46 @@ extern "C" {
#include <stdlib.h>
#include <string.h>
+#ifdef _WIN32
+#define WIN32_LEAN_AND_MEAN
+#include <windows.h>
+#else
+#include <pthread.h>
+#endif
+
+/* Sync primitives used to serialize first-time HIP runtime resolution
+ * (ensureHipLoaded) and concurrent mutations of the DynamicModules array.
+ *
+ * pthread_once / pthread_mutex on POSIX, INIT_ONCE / CRITICAL_SECTION on
+ * Windows. Kept inline in this TU to avoid a sanitizer_common dependency
+ * in the profile runtime. */
+#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);
@@ -68,12 +108,7 @@ static char DeviceArchNames[MAX_DEVICES][256];
/* Keep HIP-only to avoid an HSA dependency. */
/* -------------------------------------------------------------------------- */
-static void ensureHipLoaded(void) {
- static int Initialized = 0;
- if (Initialized)
- return;
- Initialized = 1;
-
+static void doEnsureHipLoaded(void) {
if (!__interception::DynamicLoaderAvailable()) {
if (isVerboseMode())
PROF_NOTE("%s", "Dynamic library loading not available - "
@@ -131,6 +166,21 @@ static void ensureHipLoaded(void) {
}
}
+#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 */
/* -------------------------------------------------------------------------- */
@@ -358,19 +408,23 @@ static int registerPrfSymbol(const char *Name, void *UserData) {
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 (ModuleLoadRc)
- return;
-
if (NumDynamicModules >= CapDynamicModules) {
int NewCap = CapDynamicModules ? CapDynamicModules * 2 : 64;
OffloadDynamicModuleInfo *New = (OffloadDynamicModuleInfo *)realloc(
DynamicModules, NewCap * sizeof(OffloadDynamicModuleInfo));
- if (!New)
+ if (!New) {
+ unlockDynamicModules();
return;
+ }
DynamicModules = New;
CapDynamicModules = NewCap;
}
@@ -402,9 +456,12 @@ __llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
} else if (isVerboseMode()) {
PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs);
}
+
+ unlockDynamicModules();
}
extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
+ lockDynamicModules();
for (int i = 0; i < NumDynamicModules; ++i) {
OffloadDynamicModuleInfo *MI = &DynamicModules[i];
@@ -436,11 +493,13 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
t);
}
}
+ unlockDynamicModules();
return;
}
if (isVerboseMode())
PROF_WARN("unregister called for unknown module %p\n", Ptr);
+ unlockDynamicModules();
}
/* Grow a void* array, doubling capacity (or starting at InitCap). */
@@ -806,6 +865,7 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
}
/* Dynamically-loaded modules — warn about any unprocessed TUs */
+ lockDynamicModules();
for (int i = 0; i < NumDynamicModules; ++i) {
OffloadDynamicModuleInfo *MI = &DynamicModules[i];
for (int t = 0; t < MI->NumTUs; ++t) {
@@ -816,6 +876,7 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
}
}
}
+ unlockDynamicModules();
return Ret;
}
>From 01ac776a3f23515c6371604ea416f5d1637b1209 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 19 May 2026 11:32:49 -0400
Subject: [PATCH 09/13] [compiler-rt][profile] Fix profraw filename collisions,
module-handle reuse, and host-buffer use-after-free in offload PGO
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Four bugs in the device profile drain path:
1. Filename collision (dynamic load).
processDeviceOffloadPrf was called with Target = ArchName for every
TU drain, so back-to-back unloads all resolved to the same
gfx<arch>.profile.<pid>.profraw and overwrote each other. Build
Target = "<arch>.<TUIndex>" so each per-TU drain writes to its own
file. Drop a dead snprintf inside processDeviceOffloadPrf that built
TUIndexStr but never used it.
2. Module-handle address reuse.
HIP recycles hipModule_t addresses after hipModuleUnload. After a
register/unload pair, a subsequent register of the recycled address
ended up in a new DynamicModules[] slot, but the linear ModulePtr
match in unregister found the old (drained) slot first and skipped
the new one. Clear MI->ModulePtr to nullptr after draining all TUs
in a slot, and skip cleared slots in the exit-time warning.
3. Host-buffer use-after-free.
processDeviceOffloadPrf caches freshly malloc'd host copies of the
device sections to deduplicate the RDC-mode multi-shadow case where
N shadow variables all point to the same combined section data. The
producer of the cache both stored the pointer and freed it via the
RAII cleanup, leaving the cache dangling for the next call. Tell the
cleanup to skip the free once ownership transfers to the cache.
4. Filename collision (static load, RDC).
__llvm_profile_hip_collect_device_data iterates NumShadowVariables
shadow vars (one per TU) and calls processShadowVariable for each
with Target = ArchName. Same symptom as bug 1 — all drains overwrite
the same profraw. Encode the shadow index when NumShadowVariables>1
(single-shadow programs keep the bare arch target unchanged).
Static-load single-module filename is unchanged; multi-load and RDC
multi-shadow paths now produce one profraw per TU.
---
.../profile/InstrProfilingPlatformROCm.cpp | 44 ++++++++++++++++---
1 file changed, 39 insertions(+), 5 deletions(-)
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index 24f10cd02e996..7dc0558462c46 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -465,6 +465,9 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
for (int i = 0; i < NumDynamicModules; ++i) {
OffloadDynamicModuleInfo *MI = &DynamicModules[i];
+ /* HIP recycles hipModule_t addresses after unload. Drained slots
+ * are cleared (ModulePtr = nullptr) below so a recycled-handle
+ * lookup finds the new slot, not the dead one. */
if (MI->ModulePtr != Ptr)
continue;
@@ -486,13 +489,22 @@ extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) {
int CurDev = 0;
hipGetDevice(&CurDev);
const char *ArchName = getDeviceArchName(CurDev);
- if (processDeviceOffloadPrf(TU->DeviceVar, TUIndex, ArchName) == 0)
+ /* Encode TUIndex in Target so each drain writes to its own
+ * profraw ("<arch>.<TUIndex>.profile.<pid>.profraw"); otherwise
+ * back-to-back drains overwrite the same file. Static-load
+ * (processShadowVariable) keeps the bare arch target. */
+ char TargetWithTU[64];
+ snprintf(TargetWithTU, sizeof(TargetWithTU), "%s.%d", ArchName,
+ TUIndex);
+ if (processDeviceOffloadPrf(TU->DeviceVar, TUIndex, TargetWithTU) == 0)
TU->Processed = 1;
else
PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr,
t);
}
}
+ /* Mark the slot dead so a recycled handle finds the new slot. */
+ MI->ModulePtr = nullptr;
unlockDynamicModules();
return;
}
@@ -694,20 +706,28 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
return -1;
}
+ /* Cache the freshly allocated host buffers so subsequent drains with the
+ * same device-side section pointers (the RDC-mode multi-shadow case) can
+ * reuse them. Once cached, ownership transfers to the cache: tell the
+ * RAII cleanup not to free, otherwise the cache holds dangling pointers
+ * for the next call. */
if (!CntsReused && CountersSize > 0) {
CachedDevCntsBegin = DevCntsBegin;
CachedHostCnts = HostCountersBegin;
CachedCntsSize = CountersSize;
+ HostCopies.CntsReused = 1;
}
if (!DataReused && DataSize > 0) {
CachedDevDataBegin = DevDataBegin;
CachedHostData = HostDataBegin;
CachedDataSize = DataSize;
+ HostCopies.DataReused = 1;
}
if (!NamesReused && NamesSize > 0) {
CachedDevNamesBegin = DevNamesBegin;
CachedHostNames = HostNamesBegin;
CachedNamesSize = NamesSize;
+ HostCopies.NamesReused = 1;
}
if (isVerboseMode())
@@ -781,8 +801,10 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
sizeof(RelocatedData[i].Values));
}
- char TUIndexStr[16];
- snprintf(TUIndexStr, sizeof(TUIndexStr), "%d", TUIndex);
+ /* TUIndex is encoded into Target by the dynamic-load caller; the
+ * static-load path passes bare arch and would need the same treatment
+ * to support multi-shadow drains. */
+ (void)TUIndex;
ret = __llvm_write_custom_profile(
Target, (__llvm_profile_data *)BufDataBegin,
@@ -855,7 +877,17 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
PROF_NOTE("Collecting static profile data from device %d (%s)\n", Dev,
ArchName);
for (int i = 0; i < NumShadowVariables; ++i) {
- if (processShadowVariable(OffloadShadowVariables[i], i, ArchName) != 0)
+ /* Encode the shadow-variable index in Target so per-TU drains
+ * in RDC mode (multiple shadows registered into one device
+ * image) write to distinct profraw files. Single-TU programs
+ * still get the bare arch target via the i==0 case. */
+ const char *Target = ArchName;
+ char TargetWithIdx[64];
+ if (NumShadowVariables > 1) {
+ snprintf(TargetWithIdx, sizeof(TargetWithIdx), "%s.%d", ArchName, i);
+ Target = TargetWithIdx;
+ }
+ if (processShadowVariable(OffloadShadowVariables[i], i, Target) != 0)
Ret = -1;
}
}
@@ -864,10 +896,12 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
hipSetDevice(OrigDevice);
}
- /* Dynamically-loaded modules — warn about any unprocessed TUs */
+ /* Warn about unprocessed TUs; skip cleared slots (already drained). */
lockDynamicModules();
for (int i = 0; i < NumDynamicModules; ++i) {
OffloadDynamicModuleInfo *MI = &DynamicModules[i];
+ if (!MI->ModulePtr)
+ continue;
for (int t = 0; t < MI->NumTUs; ++t) {
if (!MI->TUs[t].Processed) {
PROF_WARN("dynamic module %p TU %d was not processed before exit\n",
>From 090d850b0518bac3c84ee3cf2a098e4124e2e0b8 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 19 May 2026 12:29:57 -0400
Subject: [PATCH 10/13] [compiler-rt][profile][PGO] Move HIP dynamic-module
instrumentation to compiler-rt interceptors
Move hipModuleLoad* / hipModuleUnload handling out of InstrProfiling.cpp
and into compiler-rt. The IR pass no longer matches HIP function names
or rewrites user calls.
InstrProfilingPlatformROCm.cpp (Linux only) installs interceptors for
hipModuleLoad, hipModuleLoadData, hipModuleLoadDataEx, and
hipModuleUnload via a file-scope static initializer. Each wrapper
calls the real HIP function then forwards to the existing
__llvm_profile_offload_{register,unregister}_dynamic_module entry
points.
Interception catches every call site, including ones in libraries the
IR pass cannot see. RTInterception is already linked into
clang_rt.profile, so no build-system changes are needed.
Windows is not yet supported; it would need import-table or trampoline
patching via interception_win.cpp instead of dlsym(RTLD_NEXT).
---
.../profile/InstrProfilingPlatformROCm.cpp | 55 +++++++++
.../Instrumentation/InstrProfiling.cpp | 108 +-----------------
2 files changed, 59 insertions(+), 104 deletions(-)
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index 7dc0558462c46..7e4ce395d1bb1 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -914,3 +914,58 @@ extern "C" int __llvm_profile_hip_collect_device_data(void) {
return Ret;
}
+
+/* Interceptors for hipModuleLoad* / hipModuleUnload.
+ *
+ * Linux only. Windows would need import-table or trampoline patching
+ * via interception_win.cpp instead of dlsym(RTLD_NEXT). */
+
+#if defined(__linux__) && !defined(_WIN32)
+
+INTERCEPTOR(int, hipModuleLoad, void **module, const char *fname) {
+ int rc = REAL(hipModuleLoad)(module, fname);
+ /* Pass NULL image: no in-memory ELF is available for filename loads,
+ * so the register hook skips symbol enumeration. */
+ __llvm_profile_offload_register_dynamic_module(rc, module, nullptr);
+ return rc;
+}
+
+INTERCEPTOR(int, hipModuleLoadData, void **module, const void *image) {
+ int rc = REAL(hipModuleLoadData)(module, image);
+ __llvm_profile_offload_register_dynamic_module(rc, module, image);
+ return rc;
+}
+
+INTERCEPTOR(int, hipModuleLoadDataEx, void **module, const void *image,
+ unsigned numOptions, void **options, void **optionValues) {
+ int rc = REAL(hipModuleLoadDataEx)(module, image, numOptions, options,
+ optionValues);
+ __llvm_profile_offload_register_dynamic_module(rc, module, image);
+ return rc;
+}
+
+INTERCEPTOR(int, hipModuleUnload, void *module) {
+ /* Drain counters before the module is destroyed; device addresses
+ * captured at register time are invalid after unload. */
+ __llvm_profile_offload_unregister_dynamic_module(module);
+ return REAL(hipModuleUnload)(module);
+}
+
+/* Runs at C++ dynamic init time, before any user hipModuleLoad* call.
+ *
+ * INTERCEPT_FUNCTION must run unconditionally: our wrapper symbol
+ * preempts libamdhip64.so's at link time, so REAL() must be populated
+ * or the first user call segfaults. */
+static int installHipModuleInterceptors() {
+ if (isVerboseMode())
+ PROF_NOTE("%s", "Installing hipModuleLoad*/hipModuleUnload interceptors\n");
+ INTERCEPT_FUNCTION(hipModuleLoad);
+ INTERCEPT_FUNCTION(hipModuleLoadData);
+ INTERCEPT_FUNCTION(hipModuleLoadDataEx);
+ INTERCEPT_FUNCTION(hipModuleUnload);
+ return 0;
+}
+
+static int HipModuleInterceptorsInstalled = installHipModuleInterceptors();
+
+#endif /* __linux__ */
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index 5bb205ed3c46a..ed99161fac28c 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -447,12 +447,6 @@ class InstrLowerer final {
/// Create HIP device variable registration for profile symbols
void createHIPDeviceVariableRegistration();
-
- /// Create HIP dynamic module registration call
- void createHIPDynamicModuleRegistration();
-
- /// Create HIP dynamic module unregistration call
- void createHIPDynamicModuleUnregistration();
};
///
@@ -1041,8 +1035,10 @@ bool InstrLowerer::lower() {
// symbols
createHIPDeviceVariableRegistration();
- createHIPDynamicModuleRegistration();
- createHIPDynamicModuleUnregistration();
+ // HIP dynamic-module instrumentation (hipModuleLoad* / hipModuleUnload)
+ // is handled by interceptors in
+ // compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp — no IR-level
+ // rewrite is performed here.
// Emit runtime hook for the cases where the target does not unconditionally
// require pulling in profile runtime, and coverage is enabled on code that is
@@ -2064,102 +2060,6 @@ void InstrLowerer::emitVNodes() {
UsedVars.push_back(VNodesVar);
}
-void InstrLowerer::createHIPDynamicModuleRegistration() {
- if (isGPUProfTarget(M))
- return;
- StringRef FuncNames[] = {"hipModuleLoad", "hipModuleLoadData",
- "hipModuleLoadDataEx"};
- for (StringRef FuncName : FuncNames) {
- Function *F = M.getFunction(FuncName);
- if (!F)
- continue;
-
- for (User *U : F->users()) {
- if (auto *CB = dyn_cast<CallBase>(U)) {
- Instruction *InsertPt = nullptr;
- // If the call is an invoke instruction, we should insert the
- // registration call in the normal destination block.
- if (auto *Invoke = dyn_cast<InvokeInst>(CB)) {
- InsertPt = &*Invoke->getNormalDest()->getFirstInsertionPt();
- } else if (CB->isTerminator()) {
- // If it's another kind of terminator (e.g., callbr), we don't
- // know the semantics of the successors, so we conservatively
- // skip it. The hipModuleLoad* functions are not expected to be
- // used in other terminator instructions.
- continue;
- } else {
- // This is a normal call instruction, so we can insert after it.
- InsertPt = CB->getNextNode();
- }
-
- // If there's no valid insertion point (e.g., a malformed block),
- // skip.
- if (!InsertPt)
- continue;
-
- IRBuilder<> Builder(InsertPt);
- auto *VoidTy = Type::getVoidTy(M.getContext());
- auto *VoidPtrTy = PointerType::getUnqual(M.getContext());
- auto *Int32Ty = Type::getInt32Ty(M.getContext());
- // register(int rc, void **modulePtr, const void *image)
- auto *RegisterDynamicModuleTy =
- FunctionType::get(VoidTy, {Int32Ty, VoidPtrTy, VoidPtrTy}, false);
- FunctionCallee RegisterFunc = M.getOrInsertFunction(
- "__llvm_profile_offload_register_dynamic_module",
- RegisterDynamicModuleTy);
-
- // Arg 0: return value of the hipModuleLoad* call (hipError_t / i32).
- Value *ReturnValue = CB;
- // Arg 1: module handle (out-parameter, hipModule_t*).
- Value *ModuleHandle = CB->getArgOperand(0);
- // Arg 2: code object image pointer.
- // For hipModuleLoadData(module, image) and
- // hipModuleLoadDataEx(module, image, ...), image is arg 1.
- // For hipModuleLoad(module, fname), arg 1 is a filename — pass NULL.
- Value *ImagePtr;
- if (FuncName == "hipModuleLoad")
- ImagePtr =
- ConstantPointerNull::get(PointerType::getUnqual(M.getContext()));
- else
- ImagePtr = CB->getArgOperand(1);
-
- Builder.CreateCall(RegisterFunc, {ReturnValue, ModuleHandle, ImagePtr});
- }
- }
- }
-}
-
-void InstrLowerer::createHIPDynamicModuleUnregistration() {
- Function *F = M.getFunction("hipModuleUnload");
- if (!F)
- return;
-
- for (User *U : F->users()) {
- if (auto *CB = dyn_cast_or_null<CallBase>(U)) {
- // The insertion point is right before the call to hipModuleUnload.
- Instruction *InsertPt = CB;
-
- IRBuilder<> Builder(InsertPt);
- auto *VoidTy = Type::getVoidTy(M.getContext());
- auto *VoidPtrTy = PointerType::getUnqual(M.getContext());
-
- auto *UnregisterDynamicModuleTy =
- FunctionType::get(VoidTy, {VoidPtrTy}, false);
- FunctionCallee UnregisterFunc = M.getOrInsertFunction(
- "__llvm_profile_offload_unregister_dynamic_module",
- UnregisterDynamicModuleTy);
-
- // The argument is the module handle, which is the first
- // argument to the hipModuleUnload call.
- Value *ModuleHandle = CB->getArgOperand(0);
- Value *CastedModuleHandle =
- Builder.CreatePointerCast(ModuleHandle, VoidPtrTy);
-
- Builder.CreateCall(UnregisterFunc, {CastedModuleHandle});
- }
- }
-}
-
void InstrLowerer::emitNameData() {
if (ReferencedNames.empty())
return;
>From c5b2595909ec025703dd08b90b71a524641c5171 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 19 May 2026 14:17:12 -0400
Subject: [PATCH 11/13] [clang][PGO] Move HIP offload-PGO sections-table
emission to CGCUDANV
Emit the per-TU __llvm_profile_sections_<CUID> global from Clang
instead of the InstrProfiling pass. The middle-end pass no longer
emits HIP-runtime registration glue or reads HIP-specific globals.
clang/lib/CodeGen/CGCUDANV.cpp:
- New emitOffloadProfilingSections() called from finalizeModule() on
both device and host compiles when LangOpts.HIP and PGO are on.
- Device side: emits __llvm_profile_sections_<CUID> as a populated
7-pointer struct holding the __start_/__stop_ bounds of the
__llvm_prf_{names,cnts,data} sections plus __llvm_profile_raw_version.
ELF linker-defined start/stop symbols fill in the addresses.
- Host side: emits a void* placeholder shadow of the same name, stored
in OffloadProfShadow.
- makeRegisterGlobalsFn() (non-RDC path) and createOffloadingEntries()
(RDC path) add a __hipRegisterVar call and a small ctor that calls
__llvm_profile_offload_register_shadow_variable for the shadow, so the
host runtime can locate the device global and drain its counters.
llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp:
- Remove createHIPDeviceVariableRegistration (~115 lines) and
createProfileSectionSymbols (~35 lines): emission is now in Clang.
- Remove getCUIDFromModule, cacheGPUCUID, and the CachedCUID field.
- Drop the __hip_cuid_ scan in emitNameData (the per-TU names global
uses PrivateLinkage and is internalized at link time, so a CUID
suffix is not needed to avoid RDC link collisions).
compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp:
- The per-TU global is now the populated 7-pointer struct itself, not
a pointer indirection to __llvm_profile_sections. Drop one hipMemcpy
per drain in processShadowVariable and registerPrfSymbol.
- Rename the symbol-prefix matches from __llvm_offload_prf_ to
__llvm_profile_sections_.
---
clang/lib/CodeGen/CGCUDANV.cpp | 153 +++++++++++++
.../profile/InstrProfilingPlatformROCm.cpp | 50 ++---
.../Instrumentation/InstrProfiling.cpp | 208 ------------------
3 files changed, 175 insertions(+), 236 deletions(-)
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 3eda4237b0549..a41cfe8875837 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -28,6 +28,7 @@
#include "llvm/IR/ReplaceConstant.h"
#include "llvm/Support/Format.h"
#include "llvm/Support/VirtualFileSystem.h"
+#include "llvm/Transforms/Utils/ModuleUtils.h"
using namespace clang;
using namespace CodeGen;
@@ -72,6 +73,11 @@ class CGNVCUDARuntime : public CGCUDARuntime {
/// ModuleCtorFunction() and used to create corresponding cleanup calls in
/// ModuleDtorFunction()
llvm::GlobalVariable *GpuBinaryHandle = nullptr;
+ /// Host-side shadow for the per-TU __llvm_profile_sections_<CUID> global,
+ /// emitted only for HIP host compiles when PGO is on. Registered via
+ /// __hipRegisterVar (non-RDC) or an offloading entry (RDC) so the runtime
+ /// can locate the device-side table by name.
+ llvm::GlobalVariable *OffloadProfShadow = nullptr;
/// Whether we generate relocatable device code.
bool RelocatableDeviceCode;
/// Mangle context for device.
@@ -176,6 +182,13 @@ class CGNVCUDARuntime : public CGCUDARuntime {
void transformManagedVars();
/// Create offloading entries to register globals in RDC mode.
void createOffloadingEntries();
+ /// For HIP+PGO, emit the per-TU __llvm_profile_sections_<CUID> global.
+ /// On the device side it is the populated 7-pointer section-bounds table.
+ /// On the host side it is a placeholder void* shadow stored in
+ /// OffloadProfShadow, registered later by makeRegisterGlobalsFn (non-RDC)
+ /// or createOffloadingEntries (RDC) so the runtime can locate the
+ /// device-side table by name.
+ void emitOffloadProfilingSections();
public:
CGNVCUDARuntime(CodeGenModule &CGM);
@@ -735,6 +748,33 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
}
}
+ // Register the per-TU offload-profiling shadow so the host runtime can
+ // locate the matching device-side __llvm_profile_sections_<CUID>. We
+ // emit both __hipRegisterVar (so the HIP runtime can map the host
+ // shadow to the device symbol) and
+ // __llvm_profile_offload_register_shadow_variable (so the profile
+ // runtime adds the shadow to its drain list).
+ if (OffloadProfShadow) {
+ llvm::Constant *Name =
+ makeConstantString(std::string(OffloadProfShadow->getName()));
+ llvm::Value *RegisterVarArgs[] = {
+ &GpuBinaryHandlePtr,
+ OffloadProfShadow,
+ Name,
+ Name,
+ llvm::ConstantInt::get(IntTy, /*Extern=*/0),
+ llvm::ConstantInt::get(VarSizeTy,
+ CGM.getDataLayout().getPointerSize()),
+ llvm::ConstantInt::get(IntTy, /*Constant=*/0),
+ llvm::ConstantInt::get(IntTy, 0)};
+ Builder.CreateCall(RegisterVar, RegisterVarArgs);
+
+ llvm::FunctionCallee RegisterShadow = CGM.CreateRuntimeFunction(
+ llvm::FunctionType::get(VoidTy, {PtrTy}, false),
+ "__llvm_profile_offload_register_shadow_variable");
+ Builder.CreateCall(RegisterShadow, {OffloadProfShadow});
+ }
+
Builder.CreateRetVoid();
return RegisterKernelsFunc;
}
@@ -1256,11 +1296,124 @@ void CGNVCUDARuntime::createOffloadingEntries() {
I.Flags.getSurfTexType());
}
}
+
+ // Register the per-TU offload-profiling shadow. The offloading entry
+ // makes the linker-wrapper emit the host __hipRegisterVar call in the
+ // combined ctor. Separately emit a per-TU ctor that registers the
+ // shadow with the profile runtime's drain list.
+ if (OffloadProfShadow) {
+ llvm::offloading::emitOffloadingEntry(
+ M, Kind, OffloadProfShadow, OffloadProfShadow->getName(),
+ CGM.getDataLayout().getPointerSize(),
+ llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
+
+ llvm::LLVMContext &Ctx = M.getContext();
+ auto *PtrTy = llvm::PointerType::getUnqual(Ctx);
+ llvm::FunctionCallee RegisterShadow = CGM.CreateRuntimeFunction(
+ llvm::FunctionType::get(VoidTy, {PtrTy}, false),
+ "__llvm_profile_offload_register_shadow_variable");
+ auto *CtorFn = llvm::Function::Create(
+ llvm::FunctionType::get(VoidTy, false),
+ llvm::GlobalValue::InternalLinkage,
+ "__llvm_profile_register_shadow." + CGM.getContext().getCUIDHash(), &M);
+ auto *Entry = llvm::BasicBlock::Create(Ctx, "entry", CtorFn);
+ llvm::IRBuilder<> B(Entry);
+ B.CreateCall(RegisterShadow, {OffloadProfShadow});
+ B.CreateRetVoid();
+ llvm::appendToGlobalCtors(M, CtorFn, /*Priority=*/65535);
+ }
+}
+
+// For HIP host+device compiles with PGO enabled, emit the per-TU global
+// __llvm_profile_sections_<CUID>. Device side: a 7-pointer struct holding
+// section start/stop bounds for the names/counters/data sections plus the
+// raw-version variable. Host side: an opaque void* shadow whose only
+// purpose is to give the host-runtime a registered symbol name to look up
+// via hipGetSymbolAddress; the actual device-side data lives in the
+// matching device-side global.
+void CGNVCUDARuntime::emitOffloadProfilingSections() {
+ if (!CGM.getLangOpts().HIP)
+ return;
+ if (!CGM.getCodeGenOpts().hasProfileInstr())
+ return;
+
+ StringRef CUIDHash = CGM.getContext().getCUIDHash();
+ if (CUIDHash.empty())
+ return;
+
+ llvm::Module &M = CGM.getModule();
+ llvm::LLVMContext &Ctx = M.getContext();
+ std::string Name = ("__llvm_profile_sections_" + CUIDHash).str();
+
+ // If the global already exists (e.g. another TU was merged in), don't
+ // duplicate it.
+ if (M.getNamedValue(Name))
+ return;
+
+ if (CGM.getLangOpts().CUDAIsDevice) {
+ // Device side: emit the populated struct. Section start/stop symbols
+ // are linker-defined (ELF auto-generates __start_/__stop_ for any
+ // section whose name is a valid C identifier; AMDGPU is ELF).
+ unsigned GlobalAS = M.getDataLayout().getDefaultGlobalsAddressSpace();
+ auto *PtrTy = llvm::PointerType::get(Ctx, GlobalAS);
+ auto getOrDeclare = [&](StringRef SymName) {
+ if (auto *GV = M.getNamedGlobal(SymName))
+ return GV;
+ auto *GV = new llvm::GlobalVariable(
+ M, llvm::Type::getInt8Ty(Ctx), /*isConstant=*/false,
+ llvm::GlobalValue::ExternalLinkage, /*Initializer=*/nullptr, SymName,
+ /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
+ GlobalAS);
+ GV->setVisibility(llvm::GlobalValue::HiddenVisibility);
+ return GV;
+ };
+ auto *VersionGV = M.getNamedGlobal("__llvm_profile_raw_version");
+ if (!VersionGV) {
+ VersionGV = new llvm::GlobalVariable(
+ M, llvm::Type::getInt64Ty(Ctx), /*isConstant=*/true,
+ llvm::GlobalValue::ExternalLinkage, /*Initializer=*/nullptr,
+ "__llvm_profile_raw_version",
+ /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
+ GlobalAS);
+ }
+
+ auto *StructTy = llvm::StructType::get(
+ Ctx, {PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy, PtrTy});
+ llvm::Constant *Fields[] = {
+ getOrDeclare("__start___llvm_prf_names"),
+ getOrDeclare("__stop___llvm_prf_names"),
+ getOrDeclare("__start___llvm_prf_cnts"),
+ getOrDeclare("__stop___llvm_prf_cnts"),
+ getOrDeclare("__start___llvm_prf_data"),
+ getOrDeclare("__stop___llvm_prf_data"),
+ VersionGV,
+ };
+ auto *Init = llvm::ConstantStruct::get(StructTy, Fields);
+ auto *GV = new llvm::GlobalVariable(
+ M, StructTy, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
+ Init, Name, /*InsertBefore=*/nullptr,
+ llvm::GlobalValue::NotThreadLocal, GlobalAS);
+ GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
+ CGM.addCompilerUsedGlobal(GV);
+ return;
+ }
+
+ // Host side: emit an opaque void* shadow. Layout doesn't matter — the
+ // runtime locates it by name via hipGetSymbolAddress and treats it as
+ // the address of the device-side struct. Registration with the HIP
+ // runtime is added by makeRegisterGlobalsFn (non-RDC) or
+ // createOffloadingEntries (RDC).
+ auto *PtrTy = llvm::PointerType::getUnqual(Ctx);
+ OffloadProfShadow = new llvm::GlobalVariable(
+ M, PtrTy, /*isConstant=*/false, llvm::GlobalValue::ExternalLinkage,
+ llvm::ConstantPointerNull::get(PtrTy), Name);
+ CGM.addCompilerUsedGlobal(OffloadProfShadow);
}
// Returns module constructor to be added.
llvm::Function *CGNVCUDARuntime::finalizeModule() {
transformManagedVars();
+ emitOffloadProfilingSections();
if (CGM.getLangOpts().CUDAIsDevice) {
// Mark ODR-used device variables as compiler used to prevent it from being
// eliminated by optimization. This is necessary for device variables
diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
index 7e4ce395d1bb1..b8334094ee56e 100644
--- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
+++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp
@@ -233,7 +233,7 @@ static const char *getDeviceArchName(int DeviceId) {
/* 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_offload_prf_<CUID> */
+ void *DeviceVar; /* device address of __llvm_profile_sections_<CUID> */
int Processed; /* 0 = not yet collected, 1 = data already copied */
} OffloadDynamicTUInfo;
@@ -357,26 +357,24 @@ typedef struct {
OffloadDynamicModuleInfo *ModInfo;
} EnumState;
-/* Grow the TU array inside a module entry and register one __llvm_offload_prf_*
- * symbol. Also pre-registers the corresponding per-TU section symbols with CLR
- * (needed so hipMemcpy can copy from those device addresses later). */
+/* Grow the TU array inside a module entry and register one
+ * __llvm_profile_sections_<CUID> symbol. Also pre-registers the corresponding
+ * per-TU section symbols with CLR (needed so hipMemcpy can copy from those
+ * device addresses later). */
static int registerPrfSymbol(const char *Name, void *UserData) {
EnumState *S = (EnumState *)UserData;
OffloadDynamicModuleInfo *MI = S->ModInfo;
- /* Look up the per-TU pointer variable, then dereference to get the
- * address of __llvm_profile_sections. */
- void *DevicePtrVar = nullptr;
+ /* Look up the device address of the per-TU __llvm_profile_sections_<CUID>
+ * struct. In the current design Clang emits the struct itself at this
+ * symbol (not a pointer indirection), so the address from hipModuleGetGlobal
+ * is what processDeviceOffloadPrf needs to hipMemcpy from. */
+ void *DeviceVar = nullptr;
size_t Bytes = 0;
- if (hipModuleGetGlobal(&DevicePtrVar, &Bytes, S->Module, Name) != 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 */
}
- void *DeviceVar = nullptr;
- if (hipMemcpy(&DeviceVar, DevicePtrVar, sizeof(void *), 2 /*DToH*/) != 0) {
- PROF_WARN("failed to read sections pointer for %s\n", Name);
- return 0;
- }
/* Grow TU array if needed. */
if (MI->NumTUs >= MI->CapTUs) {
@@ -435,7 +433,7 @@ __llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
MI->NumTUs = 0;
MI->CapTUs = 0;
- /* Enumerate all __llvm_offload_prf_<CUID> symbols in the ELF image.
+ /* Enumerate all __llvm_profile_sections_<CUID> symbols in the ELF image.
* For each one, look it up via hipModuleGetGlobal (which also registers
* the device address with CLR for later hipMemcpy) and store the entry.
*
@@ -443,7 +441,8 @@ __llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
* profiling is not yet supported. */
#if __has_include(<elf.h>)
EnumState State = {*Ptr, MI};
- enumerateElfSymbols(Image, "__llvm_offload_prf_", registerPrfSymbol, &State);
+ enumerateElfSymbols(Image, "__llvm_profile_sections_", registerPrfSymbol,
+ &State);
#else
(void)Image;
if (isVerboseMode())
@@ -452,7 +451,8 @@ __llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr,
#endif
if (MI->NumTUs == 0) {
- PROF_WARN("no __llvm_offload_prf_* symbols found in module %p\n", *Ptr);
+ PROF_WARN("no __llvm_profile_sections_* symbols found in module %p\n",
+ *Ptr);
} else if (isVerboseMode()) {
PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs);
}
@@ -823,22 +823,16 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex,
static int processShadowVariable(void *ShadowVar, int TUIndex,
const char *Target) {
- void *DevicePtrVar = nullptr;
- if (hipGetSymbolAddress(&DevicePtrVar, ShadowVar) != 0) {
+ void *DeviceSections = nullptr;
+ if (hipGetSymbolAddress(&DeviceSections, ShadowVar) != 0) {
PROF_WARN("failed to get symbol address for shadow variable %p\n",
ShadowVar);
return -1;
}
- // The shadow variable is a pointer to __llvm_profile_sections (defined
- // in the GPU profile runtime). Dereference to get the struct address.
- void *DeviceOffloadPrf = nullptr;
- if (hipMemcpy(&DeviceOffloadPrf, DevicePtrVar, sizeof(void *), 2 /*DToH*/) !=
- 0) {
- PROF_WARN("failed to read sections pointer from shadow variable %p\n",
- ShadowVar);
- return -1;
- }
- return processDeviceOffloadPrf(DeviceOffloadPrf, TUIndex, Target);
+ /* The shadow's device address is the address of the per-TU
+ * __llvm_profile_sections_<CUID> struct itself. processDeviceOffloadPrf
+ * reads the 7 pointers from there with one hipMemcpy. */
+ return processDeviceOffloadPrf(DeviceSections, TUIndex, Target);
}
/* Check if HIP runtime is available and loaded */
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index ed99161fac28c..fc7adef6734d0 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -246,22 +246,6 @@ static bool profDataReferencedByCode(const Module &M) {
return enablesValueProfiling(M);
}
-// Extract CUID (Compilation Unit ID) from the module.
-// HIP/CUDA modules have a global variable __hip_cuid_<hash> that uniquely
-// identifies each translation unit. Returns an empty StringRef if not found.
-// The returned StringRef points into the GlobalVariable's name, which is
-// owned by the Module and stable for the Module's lifetime.
-static StringRef getCUIDFromModule(const Module &M) {
- for (const GlobalVariable &GV : M.globals()) {
- if (!GV.hasExternalLinkage())
- continue;
- StringRef Name = GV.getName();
- if (Name.consume_front("__hip_cuid_"))
- return Name;
- }
- return "";
-}
-
class InstrLowerer final {
public:
InstrLowerer(Module &M, const InstrProfOptions &Options,
@@ -309,7 +293,6 @@ class InstrLowerer final {
size_t NamesSize = 0;
StructType *ProfileDataTy = nullptr;
- std::string CachedCUID; // CUID cached for consistent section naming
// vector of counter load/store pairs to be register promoted.
std::vector<LoadStorePair> PromotionCandidates;
@@ -435,18 +418,8 @@ class InstrLowerer final {
/// and for any profile output file that was specified.
void emitInitialization();
- /// For GPU targets: cache the CUID for consistent section naming.
- void cacheGPUCUID();
-
/// Return the __llvm_profile_data struct type.
StructType *getProfileDataTy();
-
- /// Create __llvm_offload_prf structure for GPU targets.
- /// All sections use linker-defined __start_/__stop_ bounds.
- void createProfileSectionSymbols();
-
- /// Create HIP device variable registration for profile symbols
- void createHIPDeviceVariableRegistration();
};
///
@@ -978,8 +951,6 @@ bool InstrLowerer::lower() {
if (!ContainsProfiling && !CoverageNamesVar)
return MadeChange;
- cacheGPUCUID();
-
// We did not know how many value sites there would be inside
// the instrumented function. This is counting the number of instrumented
// target value sites to enter it as field in the profile data variable.
@@ -1028,18 +999,6 @@ bool InstrLowerer::lower() {
emitNameData();
emitVTableNames();
- // Create start/stop symbols for device code profile sections
- createProfileSectionSymbols();
-
- // Create host shadow variables and registration calls for HIP device profile
- // symbols
- createHIPDeviceVariableRegistration();
-
- // HIP dynamic-module instrumentation (hipModuleLoad* / hipModuleUnload)
- // is handled by interceptors in
- // compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp — no IR-level
- // rewrite is performed here.
-
// Emit runtime hook for the cases where the target does not unconditionally
// require pulling in profile runtime, and coverage is enabled on code that is
// not eliminated by the front-end, e.g. unused functions with internal
@@ -2074,12 +2033,6 @@ void InstrLowerer::emitNameData() {
auto *NamesVal =
ConstantDataArray::getString(Ctx, StringRef(CompressedNameStr), false);
std::string NamesVarName = std::string(getInstrProfNamesVarName());
- if (isGPUProfTarget(M)) {
- StringRef CUID =
- CachedCUID.empty() ? getCUIDFromModule(M) : StringRef(CachedCUID);
- if (!CUID.empty())
- NamesVarName = (Twine(NamesVarName) + "_" + CUID).str();
- }
NamesVar =
new GlobalVariable(M, NamesVal->getType(), true,
GlobalValue::PrivateLinkage, NamesVal, NamesVarName);
@@ -2321,165 +2274,4 @@ StructType *InstrLowerer::getProfileDataTy() {
return ProfileDataTy;
}
-void InstrLowerer::cacheGPUCUID() {
- if (!isGPUProfTarget(M))
- return;
- CachedCUID = getCUIDFromModule(M);
-}
-
-// Create CUID-suffixed pointer to __llvm_profile_sections for GPU targets.
-// The basic HIP/offload runtime exposes the original 7-entry section table:
-// names_start, names_stop, cnts_start, cnts_stop, data_start, data_stop,
-// raw_version.
-// We create a per-TU global that points to it, giving the host a unique
-// symbol for shadow variable registration.
-void InstrLowerer::createProfileSectionSymbols() {
- if (!isGPUProfTarget(M) || CachedCUID.empty())
- return;
-
- auto &Ctx = M.getContext();
- unsigned AS = M.getDataLayout().getDefaultGlobalsAddressSpace();
- auto *Int8PtrTy = PointerType::get(Ctx, AS);
-
- // __llvm_profile_sections is an array of 7 pointers defined in the GPU
- // profile runtime (InstrProfilingPlatformGPU.c). Declare it as external.
- auto *SectionsTy = ArrayType::get(Int8PtrTy, 7);
- auto *SectionsGV = M.getGlobalVariable("__llvm_profile_sections");
- if (!SectionsGV) {
- SectionsGV = new GlobalVariable(M, SectionsTy, /*isConstant=*/true,
- GlobalValue::ExternalLinkage, nullptr,
- "__llvm_profile_sections", nullptr,
- GlobalValue::NotThreadLocal, AS);
- SectionsGV->setVisibility(GlobalValue::HiddenVisibility);
- }
-
- // Create a CUID-suffixed global that stores a pointer to the sections
- // struct. Aliases can't point to declarations, so we use a pointer global.
- // The host reads through this indirection: hipGetSymbolAddress gives the
- // pointer global's device address, then one DtoH copy yields the sections
- // struct address, then another DtoH copy reads the actual sections.
- auto *PtrTy = SectionsGV->getType();
- auto *PtrInit =
- ConstantExpr::getPointerBitCastOrAddrSpaceCast(SectionsGV, PtrTy);
- std::string PtrName = "__llvm_offload_prf_" + CachedCUID;
- auto *PtrGV = new GlobalVariable(
- M, PtrTy, /*isConstant=*/true, GlobalValue::ExternalLinkage, PtrInit,
- PtrName, nullptr, GlobalValue::NotThreadLocal, AS);
- PtrGV->setVisibility(GlobalValue::DefaultVisibility);
- CompilerUsedVars.push_back(PtrGV);
-}
-
-void InstrLowerer::createHIPDeviceVariableRegistration() {
- if (isGPUProfTarget(M))
- return;
-
- StringRef CUID =
- CachedCUID.empty() ? getCUIDFromModule(M) : StringRef(CachedCUID);
- if (CUID.empty())
- return;
-
- auto &Ctx = M.getContext();
- auto *VoidTy = Type::getVoidTy(Ctx);
- auto *VoidPtrTy = PointerType::getUnqual(Ctx);
-
- std::string OffloadPrfName = ("__llvm_offload_prf_" + CUID).str();
- auto *OffloadPrfShadow = new GlobalVariable(
- M, VoidPtrTy, /*isConstant=*/false, GlobalValue::ExternalLinkage,
- ConstantPointerNull::get(cast<PointerType>(VoidPtrTy)), OffloadPrfName);
- CompilerUsedVars.push_back(OffloadPrfShadow);
-
- auto *RegisterShadowTy = FunctionType::get(VoidTy, {VoidPtrTy}, false);
- FunctionCallee RegisterShadowFunc = M.getOrInsertFunction(
- "__llvm_profile_offload_register_shadow_variable", RegisterShadowTy);
-
- Function *Ctor = M.getFunction("__hip_module_ctor");
- if (!Ctor) {
- // RDC mode: no __hip_module_ctor per-TU. Emit an offloading entry so the
- // linker wrapper generates __hipRegisterVar in the final module ctor.
- llvm::offloading::emitOffloadingEntry(
- M, llvm::object::OffloadKind::OFK_HIP, OffloadPrfShadow, OffloadPrfName,
- M.getDataLayout().getPointerSize(VoidPtrTy->getPointerAddressSpace()),
- llvm::offloading::OffloadGlobalEntry, /*Data=*/0);
-
- auto *CtorFn = Function::Create(FunctionType::get(VoidTy, false),
- GlobalValue::InternalLinkage,
- "__llvm_pgo_register_" + CUID, &M);
- auto *Entry = BasicBlock::Create(Ctx, "entry", CtorFn);
- IRBuilder<> B(Entry);
- B.CreateCall(RegisterShadowFunc, {OffloadPrfShadow});
- B.CreateRetVoid();
- appendToGlobalCtors(M, CtorFn, 65535);
- return;
- }
-
- // Locate the HIP fat-binary registration call and capture its return value
- Value *Handle = nullptr;
- for (BasicBlock &BB : *Ctor)
- for (Instruction &I : BB)
- if (auto *CB = dyn_cast<CallBase>(&I))
- if (Function *Callee = CB->getCalledFunction())
- if (Callee->getName() == "__hipRegisterFatBinary") {
- Handle = &I; // call result
- break;
- }
- if (!Handle)
- return;
- GlobalVariable *FatbinHandleGV = nullptr;
- if (auto *HandleInst = dyn_cast<Instruction>(Handle))
- for (Instruction *Cur = HandleInst->getNextNode(); Cur;
- Cur = Cur->getNextNode()) {
- auto *SI = dyn_cast<StoreInst>(Cur);
- if (!SI || SI->getValueOperand() != Handle)
- continue;
- if (auto *GV = dyn_cast<GlobalVariable>(
- SI->getPointerOperand()->stripPointerCasts())) {
- FatbinHandleGV = GV;
- break;
- }
- }
-
- if (!FatbinHandleGV) {
- LLVM_DEBUG(dbgs() << "store of __hipRegisterFatBinary call not found\n");
- }
-
- // Insert the new registration just before the ctor’s return
- ReturnInst *RetInst = nullptr;
- for (auto &BB : llvm::reverse(*Ctor))
- if ((RetInst = dyn_cast<ReturnInst>(BB.getTerminator())))
- break;
- if (!RetInst)
- return;
- IRBuilder<> Builder(RetInst);
-
- LLVM_DEBUG(dbgs() << "Found __hip_module_ctor, registering anchors for CUID="
- << CUID << "\n");
-
- auto *Int32Ty = Type::getInt32Ty(Ctx);
- auto *Int64Ty = Type::getInt64Ty(Ctx);
- auto *RegisterVarTy =
- FunctionType::get(VoidTy,
- {VoidPtrTy, VoidPtrTy, VoidPtrTy, VoidPtrTy, Int32Ty,
- Int64Ty, Int32Ty, Int32Ty},
- false);
- FunctionCallee RegisterVarFunc =
- M.getOrInsertFunction("__hipRegisterVar", RegisterVarTy);
- Value *HipHandle =
- FatbinHandleGV ? Builder.CreateLoad(VoidPtrTy, FatbinHandleGV) : Handle;
-
- auto *NameStr = ConstantDataArray::getString(Ctx, OffloadPrfName, true);
- auto *NameGV = new GlobalVariable(M, NameStr->getType(), true,
- GlobalValue::PrivateLinkage, NameStr,
- OffloadPrfName + ".name");
-
- Builder.CreateCall(RegisterVarFunc,
- {HipHandle, OffloadPrfShadow,
- Builder.CreatePointerCast(NameGV, VoidPtrTy),
- Builder.CreatePointerCast(NameGV, VoidPtrTy),
- Builder.getInt32(0),
- Builder.getInt64(M.getDataLayout().getPointerSize()),
- Builder.getInt32(0), Builder.getInt32(0)});
-
- Builder.CreateCall(RegisterShadowFunc, {OffloadPrfShadow});
-}
-
} // namespace
>From b5e0bbede92550e53abaa287f7ffd9c8437becee Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 19 May 2026 19:30:38 -0400
Subject: [PATCH 12/13] [PGO] Remove AMDGPU-specific paths from InstrProfiling
Two AMDGPU branches in the instrumentation pass had behavior-equivalent
generic forms:
- lowerIncrementAMDGPU duplicated the NVPTX __llvm_profile_instrument_gpu
call site. Merge into the existing isGPUProfTarget(M) branch using
getCounterAddress, the RuntimeLibcalls helper for the function name,
and CreateZExtOrTrunc on the step. Drop the AMDGPU method, its
declaration, and the IntrinsicsAMDGPU.h include.
- The __profd_* preemption-visibility tweak was gated on TT.isAMDGPU().
Replace with TT.isGPU() && TT.isOSBinFormatELF(); NVPTX takes a
different RelativeCounterPtr path and is unaffected.
InstrProfiling.cpp no longer references AMDGPU by name.
---
clang/lib/CodeGen/CGCUDANV.cpp | 7 +--
.../Instrumentation/InstrProfiling.cpp | 63 +++++--------------
2 files changed, 17 insertions(+), 53 deletions(-)
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index a41cfe8875837..b74de7e064c35 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -763,8 +763,7 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
Name,
Name,
llvm::ConstantInt::get(IntTy, /*Extern=*/0),
- llvm::ConstantInt::get(VarSizeTy,
- CGM.getDataLayout().getPointerSize()),
+ llvm::ConstantInt::get(VarSizeTy, CGM.getDataLayout().getPointerSize()),
llvm::ConstantInt::get(IntTy, /*Constant=*/0),
llvm::ConstantInt::get(IntTy, 0)};
Builder.CreateCall(RegisterVar, RegisterVarArgs);
@@ -1391,8 +1390,8 @@ void CGNVCUDARuntime::emitOffloadProfilingSections() {
auto *Init = llvm::ConstantStruct::get(StructTy, Fields);
auto *GV = new llvm::GlobalVariable(
M, StructTy, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
- Init, Name, /*InsertBefore=*/nullptr,
- llvm::GlobalValue::NotThreadLocal, GlobalAS);
+ Init, Name, /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
+ GlobalAS);
GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
CGM.addCompilerUsedGlobal(GV);
return;
diff --git a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
index fc7adef6734d0..1c92a36372765 100644
--- a/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
+++ b/llvm/lib/Transforms/Instrumentation/InstrProfiling.cpp
@@ -42,7 +42,6 @@
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
-#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/RuntimeLibcalls.h"
@@ -331,9 +330,6 @@ class InstrLowerer final {
/// Replace instrprof.increment with an increment of the appropriate value.
void lowerIncrement(InstrProfIncrementInst *Inc);
- /// AMDGPU specific implementation of lowerIncrement.
- void lowerIncrementAMDGPU(InstrProfIncrementInst *Inc);
-
/// Force emitting of name vars for unused functions.
void lowerCoverageData(GlobalVariable *CoverageNamesVar);
@@ -1202,24 +1198,23 @@ void InstrLowerer::lowerTimestamp(
}
void InstrLowerer::lowerIncrement(InstrProfIncrementInst *Inc) {
- if (TT.isAMDGPU()) {
- lowerIncrementAMDGPU(Inc);
- return;
- }
auto *Addr = getCounterAddress(Inc);
-
IRBuilder<> Builder(Inc);
if (isGPUProfTarget(M)) {
- auto *I64Ty = Builder.getInt64Ty();
+ auto *Int64Ty = Builder.getInt64Ty();
auto *PtrTy = Builder.getPtrTy();
auto *CalleeTy = FunctionType::get(Type::getVoidTy(M.getContext()),
- {PtrTy, PtrTy, I64Ty}, false);
- auto Callee =
- M.getOrInsertFunction("__llvm_profile_instrument_gpu", CalleeTy);
+ {PtrTy, PtrTy, Int64Ty}, false);
+ FunctionCallee Callee =
+ M.getOrInsertFunction(RTLIB::RuntimeLibcallsInfo::getLibcallImplName(
+ RTLIB::impl___llvm_profile_instrument_gpu),
+ CalleeTy);
Value *CastAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PtrTy);
Value *Uniform =
ConstantPointerNull::get(PointerType::getUnqual(M.getContext()));
- Builder.CreateCall(Callee, {CastAddr, Uniform, Inc->getStep()});
+ Value *StepI64 =
+ Builder.CreateZExtOrTrunc(Inc->getStep(), Int64Ty, "step.i64");
+ Builder.CreateCall(Callee, {CastAddr, Uniform, StepI64});
} else if (Options.Atomic || AtomicCounterUpdateAll ||
(Inc->getIndex()->isNullValue() && AtomicFirstCounter)) {
Builder.CreateAtomicRMW(AtomicRMWInst::Add, Addr, Inc->getStep(),
@@ -1235,37 +1230,6 @@ void InstrLowerer::lowerIncrement(InstrProfIncrementInst *Inc) {
Inc->eraseFromParent();
}
-void InstrLowerer::lowerIncrementAMDGPU(InstrProfIncrementInst *Inc) {
- IRBuilder<> Builder(Inc);
- LLVMContext &Context = M.getContext();
- auto *Int64Ty = Type::getInt64Ty(Context);
-
- auto *CounterIdx = Inc->getIndex();
-
- // --- Counter address ---
- GlobalVariable *Counters = getOrCreateRegionCounters(Inc);
- Value *Indices[] = {Builder.getInt32(0), CounterIdx};
- Value *Addr = Builder.CreateInBoundsGEP(Counters->getValueType(), Counters,
- Indices, "ctr.addr");
-
- auto *PtrTy = PointerType::getUnqual(Context);
- Value *UniformAddrArg = ConstantPointerNull::get(cast<PointerType>(PtrTy));
- Value *CastAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PtrTy);
-
- Value *IncStep = Inc->getStep();
- Value *StepI64 = Builder.CreateZExtOrTrunc(IncStep, Int64Ty, "step.i64");
-
- auto *CalleeTy = FunctionType::get(Type::getVoidTy(Context),
- {PtrTy, PtrTy, Int64Ty}, false);
- FunctionCallee IncrFn =
- M.getOrInsertFunction(RTLIB::RuntimeLibcallsInfo::getLibcallImplName(
- RTLIB::impl___llvm_profile_instrument_gpu),
- CalleeTy);
- Builder.CreateCall(IncrFn, {CastAddr, UniformAddrArg, StepI64});
-
- Inc->eraseFromParent();
-}
-
void InstrLowerer::lowerCoverageData(GlobalVariable *CoverageNamesVar) {
ConstantArray *Names =
cast<ConstantArray>(CoverageNamesVar->getInitializer());
@@ -1901,10 +1865,11 @@ void InstrLowerer::createDataVariable(InstrProfCntrInstBase *Inc) {
Linkage = GlobalValue::PrivateLinkage;
Visibility = GlobalValue::DefaultVisibility;
}
- // AMDGPU objects are always ET_DYN, so non-local symbols with default
- // visibility are preemptible. The CounterPtr label difference emits a REL32
- // relocation that lld rejects against preemptible targets.
- if (TT.isAMDGPU() && !GlobalValue::isLocalLinkage(Linkage))
+ // GPU-target ELF objects are always ET_DYN, so non-local symbols with
+ // default visibility are preemptible. The CounterPtr label difference
+ // emits a REL32 relocation that lld rejects against preemptible targets.
+ if (TT.isGPU() && TT.isOSBinFormatELF() &&
+ !GlobalValue::isLocalLinkage(Linkage))
Visibility = GlobalValue::ProtectedVisibility;
auto *Data =
new GlobalVariable(M, DataTy, false, Linkage, nullptr, DataVarName);
>From 6a1985151165233509e8a143adb51c07aec5455e Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 20 May 2026 09:38:45 -0400
Subject: [PATCH 13/13] [PGO] Generalize PGOInstrumentation to isGPU(); fix lit
test and add Clang codegen test
After moving HIP offload-PGO emission out of InstrProfiling and into
CGCUDANV (c5b259590 / b5e0bbede), three follow-ups remain:
- PGOInstrumentation: replace the explicit `TT.isAMDGPU() || TT.isNVPTX()`
test in `isValueProfilingDisabled` with `Triple::isGPU()`, matching the
rest of the GPU-PGO branch in InstrProfiling.cpp. No behavior change for
AMDGPU/NVPTX; SPIR/SPIR-V GPU targets now follow the same value-profiling
policy.
- amdgpu-contiguous-counters.ll -> amdgpu-profc-arrays.ll: the test was
checking for the per-TU `__llvm_offload_prf_<CUID>` global and the
`__hip_cuid_<CUID>` input symbol that `instrprof` used to consume. Both
are gone now (emission lives in Clang's CGCUDANV), so drop those CHECKs
and the unused input. The remaining test covers what the pass still
does on AMDGPU: per-kernel `__profc_*` arrays in `__llvm_prf_cnts` and
increments lowered to `__llvm_profile_instrument_gpu` calls. Rename to
match the new scope.
- clang/test/CodeGenHIP/offload-pgo-sections.hip: new test for the
CGCUDANV emission path. Covers the device subcompile (7-pointer
addrspace(1) struct, `compiler.used` pinning), the host compile (void*
shadow, `__hipRegisterVar`, `__llvm_profile_offload_register_shadow_variable`),
and the two negative guards (no PGO, no CUID).
---
.../test/CodeGenHIP/offload-pgo-sections.hip | 50 +++++++++++++++++++
.../Instrumentation/PGOInstrumentation.cpp | 3 +-
...ous-counters.ll => amdgpu-profc-arrays.ll} | 14 ++----
3 files changed, 54 insertions(+), 13 deletions(-)
create mode 100644 clang/test/CodeGenHIP/offload-pgo-sections.hip
rename llvm/test/Instrumentation/InstrProfiling/{amdgpu-contiguous-counters.ll => amdgpu-profc-arrays.ll} (61%)
diff --git a/clang/test/CodeGenHIP/offload-pgo-sections.hip b/clang/test/CodeGenHIP/offload-pgo-sections.hip
new file mode 100644
index 0000000000000..17c6fe7b9e609
--- /dev/null
+++ b/clang/test/CodeGenHIP/offload-pgo-sections.hip
@@ -0,0 +1,50 @@
+// REQUIRES: amdgpu-registered-target
+// REQUIRES: x86-registered-target
+
+// Verify CGCUDANV emits the per-TU __llvm_profile_sections_<CUID> global
+// for HIP+PGO compilations. Device subcompile: populated 7-pointer struct
+// in addrspace(1). Host compile: void* shadow registered with the HIP
+// runtime and with the profile runtime's drain list.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
+// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \
+// RUN: | FileCheck -check-prefix=DEV %s
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -cuid=abc \
+// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \
+// RUN: | FileCheck -check-prefix=HOST %s
+
+// Guard: no PGO -> no emission.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
+// RUN: -emit-llvm -o - -x hip %s \
+// RUN: | FileCheck -check-prefix=NONE %s
+
+// Guard: no CUID -> no emission.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -fprofile-instrument=clang -emit-llvm -o - -x hip %s \
+// RUN: | FileCheck -check-prefix=NONE %s
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+__device__ int helper(int x) { return x + 1; }
+__global__ void kernel(int *p) { *p = helper(*p); }
+
+// DEV-DAG: @__start___llvm_prf_names = external hidden addrspace(1) global i8
+// DEV-DAG: @__stop___llvm_prf_names = external hidden addrspace(1) global i8
+// DEV-DAG: @__start___llvm_prf_cnts = external hidden addrspace(1) global i8
+// DEV-DAG: @__stop___llvm_prf_cnts = external hidden addrspace(1) global i8
+// DEV-DAG: @__start___llvm_prf_data = external hidden addrspace(1) global i8
+// DEV-DAG: @__stop___llvm_prf_data = external hidden addrspace(1) global i8
+// DEV-DAG: @__llvm_profile_raw_version = external addrspace(1) constant i64
+// DEV: @__llvm_profile_sections_[[CUID:[0-9a-f]+]] = protected addrspace(1) constant {{.*}}@__start___llvm_prf_names{{.*}}@__stop___llvm_prf_names{{.*}}@__start___llvm_prf_cnts{{.*}}@__stop___llvm_prf_cnts{{.*}}@__start___llvm_prf_data{{.*}}@__stop___llvm_prf_data{{.*}}@__llvm_profile_raw_version
+// DEV: @llvm.compiler.used = {{.*}}@__llvm_profile_sections_[[CUID]]
+
+// HOST: @__llvm_profile_sections_[[CUID:[0-9a-f]+]] = global ptr null
+// HOST: @llvm.compiler.used = {{.*}}@__llvm_profile_sections_[[CUID]]
+// HOST: define internal void @__hip_register_globals
+// HOST: call void @__hipRegisterVar({{.*}}@__llvm_profile_sections_[[CUID]],
+// HOST: call void @__llvm_profile_offload_register_shadow_variable(ptr @__llvm_profile_sections_[[CUID]])
+
+// NONE-NOT: __llvm_profile_sections_
+// NONE-NOT: __llvm_profile_offload_register_shadow_variable
diff --git a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
index e0bfccd7cd54d..b6d07aa821e7f 100644
--- a/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
+++ b/llvm/lib/Transforms/Instrumentation/PGOInstrumentation.cpp
@@ -379,10 +379,9 @@ class FunctionInstrumenter final {
// values. Supporting other values is relatively straight-forward - just
// another counter range within the context.
bool isValueProfilingDisabled() const {
- const Triple &TT = M.getTargetTriple();
return DisableValueProfiling ||
InstrumentationType == PGOInstrumentationType::CTXPROF ||
- TT.isAMDGPU() || TT.isNVPTX();
+ M.getTargetTriple().isGPU();
}
bool shouldInstrumentEntryBB() const {
diff --git a/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll b/llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll
similarity index 61%
rename from llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
rename to llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll
index 2252a722a91cc..eab78fb3591b1 100644
--- a/llvm/test/Instrumentation/InstrProfiling/amdgpu-contiguous-counters.ll
+++ b/llvm/test/Instrumentation/InstrProfiling/amdgpu-profc-arrays.ll
@@ -1,16 +1,12 @@
-;; Test that AMDGPU targets use contiguous counter allocation with CUID-based naming.
-;; This avoids linker reordering issues where individual __profc_* symbols could be
-;; placed in any order within the section.
+;; Per-kernel __profc_* arrays land in section __llvm_prf_cnts with one slot
+;; per counter, and counter increments lower to __llvm_profile_instrument_gpu
+;; calls whose pointer argument is a GEP into the per-kernel array.
; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -passes=instrprof < %s | FileCheck %s
-;; Simulate a module with CUID (as generated by HIP compilation)
- at __hip_cuid_abc123 = addrspace(1) global i8 0
-
@__profn_kernel1 = private constant [7 x i8] c"kernel1"
@__profn_kernel2 = private constant [7 x i8] c"kernel2"
-;; Per-kernel counter arrays: kernel1 has 2 slots, kernel2 has 1 (section "__llvm_prf_cnts")
; CHECK: @__profc_kernel1 = linkonce_odr protected addrspace(1) global [2 x i64] zeroinitializer, section "__llvm_prf_cnts"
; CHECK: @__profc_kernel2 = linkonce_odr protected addrspace(1) global [1 x i64] zeroinitializer, section "__llvm_prf_cnts"
@@ -27,8 +23,4 @@ define amdgpu_kernel void @kernel2() {
declare void @llvm.instrprof.increment(ptr, i64, i32, i32)
-;; Registration symbol points at external section table (Joseph #187136 layout)
-; CHECK: @__llvm_offload_prf_abc123 = addrspace(1) constant ptr addrspace(1) @__llvm_profile_sections
-
-;; Second counter slot uses GEP into the same [2 x i64] arrays
; CHECK: call void @__llvm_profile_instrument_gpu(ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ([2 x i64], ptr addrspace(1) @__profc_kernel1, i32 0, i32 1) to ptr), ptr null, i64 1)
More information about the cfe-commits
mailing list