[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 17 08:31:27 PDT 2026
================
@@ -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);
----------------
yxsamliu wrote:
done
https://github.com/llvm/llvm-project/pull/177665
More information about the llvm-commits
mailing list