[compiler-rt] [llvm] [PGO][AMDGPU] Add basic HIP offload PGO support (PR #177665)

[Yaxun Liu via llvm-commits]

https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/177665

>From 6624247e111801c16c68a8f3bf570cbf9eb3e1e1 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

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        |  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
+}