[compiler-rt] [llvm] [PGO][AMDGPU] Add basic HIP offload PGO support (PR #177665)
Yaxun Liu via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 10 12:09:21 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/177665
>From 32d2983bccac877ecbd30f6d36f817e0a368e937 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] [PGO][AMDGPU] Add basic HIP offload PGO support
Squash of the HIP offload PGO stack (previously commits c2a60078a2eb through
67dd30ca6511 on offload-pgo-main): device/host profile collection for HIP,
ROCm platform runtime (InstrProfilingPlatformROCm.cpp), AMDGPU
instrumentation to __llvm_profile_instrument_gpu, review-driven cleanups
(-fno-exceptions / MSVC MultiThreadedDLL for profile+interception merge,
PGOInstrumentation const Triple&, lit fixes), and Windows lit CRT alignment
for Profile-* and ASan+coverage with the profile runtime.
---
compiler-rt/lib/profile/CMakeLists.txt | 80 +-
compiler-rt/lib/profile/InstrProfilingFile.c | 6 +
.../profile/InstrProfilingPlatformROCm.cpp | 830 ++++++++++++++++++
.../TestCases/asan_and_llvm_coverage_test.cpp | 2 +-
compiler-rt/test/asan/lit.cfg.py | 20 +
compiler-rt/test/profile/CMakeLists.txt | 10 +
.../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 +
12 files changed, 1506 insertions(+), 39 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 8c196d15841a4..b6065868681f7 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)
@@ -167,14 +179,62 @@ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx")
endif()
endif()
+# AMDGPU bare-metal profile (offload PGO): InstrProfiling.c includes <string.h> and
+# <limits.h>. Without forcing the include path, the GPU runtimes build can pick up
+# the *host* glibc headers (e.g. /usr/include/string.h -> bits/libc-header-start.h),
+# which is wrong for --target=amdgcn-amd-amdhsa -ffreestanding and fails the build.
+# Additionally, LLVM-libc headers for the GPU triple are generated by hdrgen during the
+# same ninja graph; compiling profile sources too early can run before those headers
+# exist. Fix: -nostdinc and -isystem only (1) LLVM-libc's generated include tree for
+# amdgcn-amd-amdhsa under the main LLVM build dir and (2) clang's resource-dir/include
+# for compiler builtins (stddef.h, etc.). See add_dependencies below for ordering.
+if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn" AND COMPILER_RT_PROFILE_BAREMETAL)
+ # Prefer LLVM_BINARY_DIR: the runtimes external project passes it and it points at
+ # the top-level LLVM build directory (where include/amdgcn-amd-amdhsa/ is written).
+ # Standalone compiler-rt configurations may only define LLVM_CMAKE_DIR instead.
+ set(PROFILE_RT_LLVM_BUILD_DIR "")
+ if(DEFINED LLVM_BINARY_DIR AND LLVM_BINARY_DIR)
+ set(PROFILE_RT_LLVM_BUILD_DIR "${LLVM_BINARY_DIR}")
+ elseif(DEFINED LLVM_CMAKE_DIR AND LLVM_CMAKE_DIR)
+ set(PROFILE_RT_LLVM_BUILD_DIR "${LLVM_CMAKE_DIR}")
+ endif()
+ if(PROFILE_RT_LLVM_BUILD_DIR)
+ execute_process(
+ COMMAND "${CMAKE_C_COMPILER}" -print-resource-dir
+ OUTPUT_VARIABLE PROFILE_RT_CLANG_RESOURCE_DIR
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ RESULT_VARIABLE PROFILE_RT_PRINT_RES_DIR_RC)
+ if(PROFILE_RT_PRINT_RES_DIR_RC EQUAL 0 AND PROFILE_RT_CLANG_RESOURCE_DIR)
+ set(PROFILE_RT_LIBC_GPU_INC "${PROFILE_RT_LLVM_BUILD_DIR}/include/amdgcn-amd-amdhsa")
+ list(APPEND EXTRA_FLAGS "-nostdinc")
+ list(APPEND EXTRA_FLAGS "-isystem${PROFILE_RT_LIBC_GPU_INC}")
+ list(APPEND EXTRA_FLAGS "-isystem${PROFILE_RT_CLANG_RESOURCE_DIR}/include")
+ else()
+ message(WARNING "amdgcn bare-metal profile: could not run ${CMAKE_C_COMPILER} -print-resource-dir; "
+ "GPU profile runtime may pick up host libc headers.")
+ endif()
+ endif()
+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}")
@@ -191,6 +251,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}
@@ -199,8 +260,21 @@ 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}
PARENT_TARGET profile)
+ # With -nostdinc, InstrProfiling.c / InstrProfilingWriter.c need generated
+ # llvm-libc headers. Without an explicit dependency, ninja can schedule the GPU
+ # profile objects before hdrgen runs, causing missing-header failures.
+ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn" AND COMPILER_RT_PROFILE_BAREMETAL)
+ if(TARGET clang_rt.profile-amdgcn)
+ foreach(_pgo_libc_hdr_dep IN ITEMS libc.include.limits libc.include.string)
+ if(TARGET ${_pgo_libc_hdr_dep})
+ add_dependencies(clang_rt.profile-amdgcn ${_pgo_libc_hdr_dep})
+ endif()
+ endforeach()
+ endif()
+ endif()
endif()
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/compiler-rt/test/asan/TestCases/asan_and_llvm_coverage_test.cpp b/compiler-rt/test/asan/TestCases/asan_and_llvm_coverage_test.cpp
index 776684fd8045d..96d7b343324ce 100644
--- a/compiler-rt/test/asan/TestCases/asan_and_llvm_coverage_test.cpp
+++ b/compiler-rt/test/asan/TestCases/asan_and_llvm_coverage_test.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx_asan -coverage -O0 %s -o %t
+// RUN: %clangxx_asan_with_profile -coverage -O0 %s -o %t
// RUN: %env_asan_opts=check_initialization_order=1 %run %t 2>&1 | FileCheck %s
#include <stdio.h>
diff --git a/compiler-rt/test/asan/lit.cfg.py b/compiler-rt/test/asan/lit.cfg.py
index 0194c720d003b..94bfc80ef0ba5 100644
--- a/compiler-rt/test/asan/lit.cfg.py
+++ b/compiler-rt/test/asan/lit.cfg.py
@@ -99,6 +99,20 @@ def get_required_attr(config, attr_name):
clang_asan_cflags = clang_asan_static_cflags + asan_dynamic_flags
clang_asan_cxxflags = clang_asan_static_cxxflags + asan_dynamic_flags
+# PR A can build clang_rt.profile with /MD on Windows when it merges
+# RTInterception / RTSanitizerCommon. Tests that combine ASan with -coverage
+# link that profile archive too, so give them the same CRT model override that
+# check-profile already uses.
+asan_profile_flags = []
+if platform.system() == "Windows" and target_is_msvc:
+ asan_profile_flags = [
+ "-D_MT",
+ "-D_DLL",
+ "-Wl,-nodefaultlib:libcmt,-defaultlib:msvcrt,-defaultlib:oldnames",
+ ]
+clang_asan_profile_cflags = clang_asan_cflags + asan_profile_flags
+clang_asan_profile_cxxflags = clang_asan_cxxflags + asan_profile_flags
+
# Add win32-(static|dynamic)-asan features to mark tests as passing or failing
# in those modes. lit doesn't support logical feature test combinations.
if platform.system() == "Windows":
@@ -121,6 +135,12 @@ def build_invocation(compile_flags, with_lto=False):
config.substitutions.append(("%clangxx ", build_invocation(target_cxxflags)))
config.substitutions.append(("%clang_asan ", build_invocation(clang_asan_cflags)))
config.substitutions.append(("%clangxx_asan ", build_invocation(clang_asan_cxxflags)))
+config.substitutions.append(
+ ("%clang_asan_with_profile ", build_invocation(clang_asan_profile_cflags))
+)
+config.substitutions.append(
+ ("%clangxx_asan_with_profile ", build_invocation(clang_asan_profile_cxxflags))
+)
config.substitutions.append(
("%clang_asan_lto ", build_invocation(clang_asan_cflags, True))
)
diff --git a/compiler-rt/test/profile/CMakeLists.txt b/compiler-rt/test/profile/CMakeLists.txt
index a6d8a9684508d..6e66bafcfc07b 100644
--- a/compiler-rt/test/profile/CMakeLists.txt
+++ b/compiler-rt/test/profile/CMakeLists.txt
@@ -22,6 +22,16 @@ pythonize_bool(LLVM_ENABLE_CURL)
foreach(arch ${PROFILE_TEST_ARCH})
set(PROFILE_TEST_TARGET_ARCH ${arch})
get_test_cc_for_arch(${arch} PROFILE_TEST_TARGET_CC PROFILE_TEST_TARGET_CFLAGS)
+ # When lib/profile/CMakeLists.txt builds clang_rt.profile with
+ # MultiThreadedDLL (merging RTInterception / RTSanitizerCommon), Profile-*
+ # tests must link with the same CRT model as the static archive. Otherwise
+ # the just-built profile .objs use __imp_* (DLL CRT) while the test binary
+ # defaults to /MT — LNK2019 / LNK4098 on Windows premerge.
+ # Same pattern as asan/lit.cfg.py dynamic-runtime flags for *-windows-msvc.
+ if(MSVC AND COMPILER_RT_HAS_INTERCEPTION AND NOT COMPILER_RT_PROFILE_BAREMETAL)
+ string(APPEND PROFILE_TEST_TARGET_CFLAGS
+ " -D_MT -D_DLL -Wl,-nodefaultlib:libcmt,-defaultlib:msvcrt,-defaultlib:oldnames")
+ endif()
set(CONFIG_NAME Profile-${arch})
configure_lit_site_cfg(
${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
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
+}
More information about the llvm-commits
mailing list