diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index e35ff6b03feeb..dca14ee8f7d67 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9864,6 +9864,21 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX())) LinkerArgs.emplace_back("-lompdevice"); + // With PGO/coverage instrumentation, GPU device code references the + // device profile runtime (__llvm_profile_instrument_gpu and the + // __llvm_profile_sections bounds table emitted by + // InstrProfilingPlatformGPU). The offload device link does not otherwise + // pull it in, so forward the static device profile runtime to the GPU + // device linker. The archive is arch-suffixed, so pass its full path + // rather than a -l name. + if (ToolChain::needsProfileRT(Args) && + (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX())) { + std::string ProfileRT = + TC->getCompilerRT(Args, "profile", ToolChain::FT_Static); + if (TC->getVFS().exists(ProfileRT)) + LinkerArgs.emplace_back(Args.MakeArgString(ProfileRT)); + } + // For SPIR-V, pass some extra flags to `spirv-link`, the out-of-tree // SPIR-V linker. `spirv-link` isn't called in LTO mode so restrict these // flags to normal compilation. diff --git a/clang/lib/Driver/ToolChains/HIPAMD.cpp b/clang/lib/Driver/ToolChains/HIPAMD.cpp index 72beb1a901243..c58113fea2512 100644 --- a/clang/lib/Driver/ToolChains/HIPAMD.cpp +++ b/clang/lib/Driver/ToolChains/HIPAMD.cpp @@ -19,6 +19,7 @@ #include "clang/Options/Options.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/Path.h" +#include "llvm/Support/VirtualFileSystem.h" #include "llvm/TargetParser/TargetParser.h" using namespace clang::driver; @@ -152,6 +153,25 @@ void AMDGCN::Linker::constructLldCommand(Compilation &C, const JobAction &JA, LldArgs.push_back("--no-whole-archive"); + // With PGO/coverage instrumentation, instrumented device code references the + // device profile runtime (__llvm_profile_instrument_gpu and the + // __llvm_profile_sections bounds table emitted by InstrProfilingPlatformGPU). + // The new-offload-driver path injects this in LinkerWrapper::ConstructJob, but + // HIP using the traditional offload path (e.g. on Windows, which does not + // route device linking through clang-linker-wrapper) reaches the device link + // here instead. Forward the static device profile runtime to this lld device + // link so the runtime is pulled in regardless of offload-driver/host OS. The + // archive is arch-suffixed, so pass its full path rather than a -l name. + if (ToolChain::needsProfileRT(Args)) { + std::string ProfileRT = + TC.getCompilerRT(Args, "profile", ToolChain::FT_Static); + // Use the ToolChain VFS (matches the new-offload-driver path in + // Clang.cpp) so overlay/virtual filesystems used by the driver are + // honored; llvm::sys::fs bypasses them and can wrongly skip the runtime. + if (TC.getVFS().exists(ProfileRT)) + LldArgs.push_back(Args.MakeArgString(ProfileRT)); + } + const char *Lld = Args.MakeArgStringRef(getToolChain().GetProgramPath("lld")); C.addCommand(std::make_unique(JA, *this, ResponseFileSupport::None(), Lld, LldArgs, Inputs, Output)); diff --git a/compiler-rt/lib/profile/CMakeLists.txt b/compiler-rt/lib/profile/CMakeLists.txt index e933e26cb241b..51ae053c983de 100644 --- a/compiler-rt/lib/profile/CMakeLists.txt +++ b/compiler-rt/lib/profile/CMakeLists.txt @@ -94,7 +94,12 @@ if (NOT COMPILER_RT_PROFILE_BAREMETAL) InstrProfilingValue.c ) if(COMPILER_RT_BUILD_PROFILE_ROCM) - list(APPEND PROFILE_SOURCES InstrProfilingPlatformROCm.cpp) + # Linux uses HSA introspection (InstrProfilingPlatformROCm.cpp); Windows + # keeps the legacy host-shadow drain (InstrProfilingPlatformROCmWindows.cpp). + # Each is platform-guarded internally, so adding both is safe (the other + # platform compiles to an empty TU). + list(APPEND PROFILE_SOURCES InstrProfilingPlatformROCm.cpp + InstrProfilingPlatformROCmWindows.cpp) endif() endif() @@ -181,7 +186,8 @@ if(COMPILER_RT_HAS_INTERCEPTION AND NOT COMPILER_RT_PROFILE_BAREMETAL endif() if(NOT PROFILE_HAS_HIP_INTERCEPTOR) - list(REMOVE_ITEM PROFILE_SOURCES InstrProfilingPlatformROCm.cpp) + list(REMOVE_ITEM PROFILE_SOURCES InstrProfilingPlatformROCm.cpp + InstrProfilingPlatformROCmWindows.cpp) endif() # Only advertise the ROCm interceptor to InstrProfilingFile.c when its @@ -281,7 +287,8 @@ if(COMPILER_RT_BUILD_PROFILE_ROCM AND NOT COMPILER_RT_PROFILE_BAREMETAL AND TARGET RTSanitizerCommon.${COMPILER_RT_DEFAULT_TARGET_ARCH} AND TARGET RTSanitizerCommonLibc.${COMPILER_RT_DEFAULT_TARGET_ARCH}) - set(PROFILE_ROCM_SOURCES ${PROFILE_SOURCES} InstrProfilingPlatformROCm.cpp) + set(PROFILE_ROCM_SOURCES ${PROFILE_SOURCES} InstrProfilingPlatformROCm.cpp + InstrProfilingPlatformROCmWindows.cpp) # Enables the device-collection call in InstrProfilingFile.c. set(PROFILE_ROCM_FLAGS ${EXTRA_FLAGS} -DCOMPILER_RT_BUILD_PROFILE_ROCM=1) diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp index ee00c572e3a42..c72852ba8a6f2 100644 --- a/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp +++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCm.cpp @@ -5,6 +5,40 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +// +// This is the Linux/Unix device profile drain, which decouples device counter +// collection from the host by walking HSA code objects (no host-side per-TU +// shadow). Windows has no HSA runtime, so it keeps the legacy HIP host-shadow +// mechanism in InstrProfilingPlatformROCmWindows.cpp; the CMake selects exactly +// one of the two by platform. Decoupled (host-uninstrumented) collection is +// only supported here, on Linux. +// +// Device-side profile data drain for AMDGPU via HSA introspection. +// +// At process exit this walks every loaded HSA code object on every GPU agent, +// finds the device-side __llvm_profile_sections bounds table (emitted by +// InstrProfilingPlatformGPU.c), copies its counters/data/names regions back to +// the host, and writes a target-prefixed .profraw via __llvm_write_custom_profile. +// +// The drain is decoupled from the host-side profile write: it runs from an +// atexit handler registered in a constructor, so device counters are collected +// whether or not the host translation units were instrumented, and without any +// host-side per-TU shadow variable, CUID matching, or hipModuleLoad +// interception. +// +// All HSA and HIP entry points are resolved with dlopen/dlsym (via the +// interception helpers) so libclang_rt.profile still links and runs on hosts +// without ROCm installed. +// +//===----------------------------------------------------------------------===// + +// Host-only: this drains device counters from the host process. The device +// side (the __llvm_profile_sections bounds table) is emitted by +// InstrProfilingPlatformGPU.c. When this file is compiled for a GPU target as +// part of the device profile runtime build it reduces to an empty TU. It also +// reduces to an empty TU on Windows, which uses +// InstrProfilingPlatformROCmWindows.cpp instead. +#if !defined(__NVPTX__) && !defined(__AMDGPU__) && !defined(_WIN32) extern "C" { #include "InstrProfiling.h" @@ -13,520 +47,259 @@ extern "C" { } #include "interception/interception.h" -// C library headers (not etc.): clang_rt.profile is built with +// C library headers only (not etc.): clang_rt.profile is built with // -nostdinc++ and avoids the C++ standard library (see profile/CMakeLists.txt). #include #include #include #include -#ifdef _WIN32 -#define WIN32_LEAN_AND_MEAN -#include -#else -#include -#include -#endif - -/* Serialize one-time HIP loader resolution and DynamicModules mutations. - * Inline to avoid a sanitizer_common dependency. */ -#ifdef _WIN32 -static INIT_ONCE HipLoadedOnce = INIT_ONCE_STATIC_INIT; -static CRITICAL_SECTION DynamicModulesLock; -static INIT_ONCE DynamicModulesLockInit = INIT_ONCE_STATIC_INIT; -static BOOL CALLBACK initDynamicModulesLockCb(PINIT_ONCE, PVOID, PVOID *) { - InitializeCriticalSection(&DynamicModulesLock); - return TRUE; -} -static void lockDynamicModules(void) { - InitOnceExecuteOnce(&DynamicModulesLockInit, initDynamicModulesLockCb, NULL, - NULL); - EnterCriticalSection(&DynamicModulesLock); -} -static void unlockDynamicModules(void) { - LeaveCriticalSection(&DynamicModulesLock); -} -#else -static pthread_once_t HipLoadedOnce = PTHREAD_ONCE_INIT; -static pthread_mutex_t DynamicModulesLock = PTHREAD_MUTEX_INITIALIZER; -static void lockDynamicModules(void) { - pthread_mutex_lock(&DynamicModulesLock); -} -static void unlockDynamicModules(void) { - pthread_mutex_unlock(&DynamicModulesLock); -} -#endif - -static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex, - const char *Target); - -static int isVerboseMode() { - static int IsVerbose = -1; - if (IsVerbose == -1) - IsVerbose = getenv("LLVM_PROFILE_VERBOSE") != nullptr; - return IsVerbose; -} - /* -------------------------------------------------------------------------- */ -/* Dynamic loading of HIP runtime symbols */ +/* Minimal HSA type/enum stubs */ +/* */ +/* compiler-rt cannot depend on ROCm headers at build time, so mirror just */ +/* the handful of HSA declarations the drain needs. Values match */ +/* hsa/hsa.h and hsa/hsa_ven_amd_loader.h. */ /* -------------------------------------------------------------------------- */ -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); +typedef uint32_t prof_hsa_status_t; +#define PROF_HSA_STATUS_SUCCESS ((prof_hsa_status_t)0x0) +#define PROF_HSA_STATUS_INFO_BREAK ((prof_hsa_status_t)0x1) -/* Minimal hipDeviceProp_t (HIP 6.x R0600): only gcnArchName at offset 1160 - * is read. Padded to 4096 to tolerate ABI growth. */ typedef struct { - char padding[1160]; - char gcnArchName[256]; - char tail_padding[2680]; -} HipDevicePropMinimal; -typedef int (*hipGetDevicePropertiesTy)(HipDevicePropMinimal *, int); - -static hipGetSymbolAddressTy pHipGetSymbolAddress = nullptr; -static hipMemcpyTy pHipMemcpy = nullptr; -static hipModuleGetGlobalTy pHipModuleGetGlobal = nullptr; -static hipGetDeviceCountTy pHipGetDeviceCount = nullptr; -static hipGetDeviceTy pHipGetDevice = nullptr; -static hipSetDeviceTy pHipSetDevice = nullptr; -static hipGetDevicePropertiesTy pHipGetDeviceProperties = nullptr; - -static int NumDevices = 0; -/* 256 matches hipDeviceProp_t::gcnArchName, the source field width. */ -static char (*DeviceArchNames)[256] = nullptr; - -/* -------------------------------------------------------------------------- */ -/* Device-to-host copies */ -/* Keep HIP-only to avoid an HSA dependency. */ -/* -------------------------------------------------------------------------- */ - -static void doEnsureHipLoaded(void) { - if (!__interception::DynamicLoaderAvailable()) { - if (isVerboseMode()) - PROF_NOTE("%s", "Dynamic library loading not available - " - "HIP profiling disabled\n"); - return; - } - -#ifdef _WIN32 - static const char HipLibName[] = "amdhip64.dll"; -#else - static const char HipLibName[] = "libamdhip64.so"; -#endif - - void *Handle = __interception::OpenLibrary(HipLibName); - if (!Handle) - return; - - pHipGetSymbolAddress = (hipGetSymbolAddressTy)__interception::LookupSymbol( - Handle, "hipGetSymbolAddress"); - pHipMemcpy = (hipMemcpyTy)__interception::LookupSymbol(Handle, "hipMemcpy"); - pHipModuleGetGlobal = (hipModuleGetGlobalTy)__interception::LookupSymbol( - Handle, "hipModuleGetGlobal"); - pHipGetDeviceCount = (hipGetDeviceCountTy)__interception::LookupSymbol( - Handle, "hipGetDeviceCount"); - pHipGetDevice = - (hipGetDeviceTy)__interception::LookupSymbol(Handle, "hipGetDevice"); - pHipSetDevice = - (hipSetDeviceTy)__interception::LookupSymbol(Handle, "hipSetDevice"); - pHipGetDeviceProperties = - (hipGetDevicePropertiesTy)__interception::LookupSymbol( - Handle, "hipGetDevicePropertiesR0600"); - if (!pHipGetDeviceProperties) - pHipGetDeviceProperties = - (hipGetDevicePropertiesTy)__interception::LookupSymbol( - Handle, "hipGetDeviceProperties"); - - if (pHipGetDeviceCount && pHipGetDeviceProperties) { - int Count = 0; - if (pHipGetDeviceCount(&Count) == 0 && Count > 0) { - DeviceArchNames = (char (*)[256])calloc(Count, sizeof(*DeviceArchNames)); - if (!DeviceArchNames) { - PROF_ERR("%s\n", "failed to allocate device arch name table"); - return; - } - HipDevicePropMinimal Prop; - for (int i = 0; i < Count; ++i) { - __builtin_memset(&Prop, 0, sizeof(Prop)); - if (pHipGetDeviceProperties(&Prop, i) == 0) { - strncpy(DeviceArchNames[i], Prop.gcnArchName, - sizeof(DeviceArchNames[i]) - 1); - DeviceArchNames[i][sizeof(DeviceArchNames[i]) - 1] = '\0'; - if (isVerboseMode()) - PROF_NOTE("Device %d arch: %s\n", i, DeviceArchNames[i]); - } - } - NumDevices = Count; - } - } -} - -#ifdef _WIN32 -static BOOL CALLBACK ensureHipLoadedCb(PINIT_ONCE, PVOID, PVOID *) { - doEnsureHipLoaded(); - return TRUE; -} -#endif - -static void ensureHipLoaded(void) { -#ifdef _WIN32 - InitOnceExecuteOnce(&HipLoadedOnce, ensureHipLoadedCb, NULL, NULL); -#else - pthread_once(&HipLoadedOnce, doEnsureHipLoaded); -#endif -} + uint64_t handle; +} prof_hsa_agent_t; +typedef struct { + uint64_t handle; +} prof_hsa_executable_t; +typedef struct { + uint64_t handle; +} prof_hsa_executable_symbol_t; -/* -------------------------------------------------------------------------- */ -/* Public wrappers that forward to the loaded HIP symbols */ -/* -------------------------------------------------------------------------- */ +typedef uint32_t prof_hsa_agent_info_t; +#define PROF_HSA_AGENT_INFO_NAME ((prof_hsa_agent_info_t)0) +#define PROF_HSA_AGENT_INFO_DEVICE ((prof_hsa_agent_info_t)17) -static int hipGetSymbolAddress(void **devPtr, const void *symbol) { - ensureHipLoaded(); - return pHipGetSymbolAddress ? pHipGetSymbolAddress(devPtr, symbol) : -1; -} +typedef uint32_t prof_hsa_device_type_t; +#define PROF_HSA_DEVICE_TYPE_GPU ((prof_hsa_device_type_t)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; -} +typedef uint32_t prof_hsa_symbol_kind_t; +#define PROF_HSA_SYMBOL_KIND_VARIABLE ((prof_hsa_symbol_kind_t)0) -/* Device section symbols must be registered with CLR first; otherwise - * hipMemcpy may take a CPU path and crash. */ -static int memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) { - return hipMemcpy(Dst, Src, Size, 2 /* DToH */); -} +typedef uint32_t prof_hsa_executable_symbol_info_t; +#define PROF_HSA_EXECUTABLE_SYMBOL_INFO_TYPE \ + ((prof_hsa_executable_symbol_info_t)0) +#define PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH \ + ((prof_hsa_executable_symbol_info_t)1) +#define PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME \ + ((prof_hsa_executable_symbol_info_t)2) +#define PROF_HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS \ + ((prof_hsa_executable_symbol_info_t)21) -static int hipModuleGetGlobal(void **DevPtr, size_t *Bytes, void *Module, - const char *Name) { - ensureHipLoaded(); - return pHipModuleGetGlobal ? pHipModuleGetGlobal(DevPtr, Bytes, Module, Name) - : -1; -} +#define PROF_HSA_EXTENSION_AMD_LOADER ((uint16_t)0x201) -static int hipGetDevice(int *DeviceId) { - ensureHipLoaded(); - return pHipGetDevice ? pHipGetDevice(DeviceId) : -1; -} +typedef uint32_t prof_hsa_loader_storage_type_t; -static int hipSetDevice(int DeviceId) { - ensureHipLoaded(); - return pHipSetDevice ? pHipSetDevice(DeviceId) : -1; -} +typedef struct { + prof_hsa_agent_t agent; + prof_hsa_executable_t executable; + prof_hsa_loader_storage_type_t code_object_storage_type; + const void *code_object_storage_base; + size_t code_object_storage_size; + size_t code_object_storage_offset; + const void *segment_base; + size_t segment_size; +} prof_hsa_loader_segment_descriptor_t; + +typedef prof_hsa_status_t (*hsa_init_ty)(void); +typedef prof_hsa_status_t (*hsa_iterate_agents_ty)( + prof_hsa_status_t (*)(prof_hsa_agent_t, void *), void *); +typedef prof_hsa_status_t (*hsa_agent_get_info_ty)(prof_hsa_agent_t, + prof_hsa_agent_info_t, + void *); +typedef prof_hsa_status_t (*hsa_executable_iterate_agent_symbols_ty)( + prof_hsa_executable_t, prof_hsa_agent_t, + prof_hsa_status_t (*)(prof_hsa_executable_t, prof_hsa_agent_t, + prof_hsa_executable_symbol_t, void *), + void *); +typedef prof_hsa_status_t (*hsa_executable_symbol_get_info_ty)( + prof_hsa_executable_symbol_t, prof_hsa_executable_symbol_info_t, void *); +typedef prof_hsa_status_t (*hsa_system_get_major_extension_table_ty)( + uint16_t, uint16_t, size_t, void *); +typedef prof_hsa_status_t (*hsa_loader_query_segment_descriptors_ty)( + prof_hsa_loader_segment_descriptor_t *, size_t *); + +/* First two members of hsa_ven_amd_loader_1_00_pfn_t. Only + * query_segment_descriptors is used; query_host_address keeps the offset. */ +typedef struct { + void *query_host_address; + hsa_loader_query_segment_descriptors_ty query_segment_descriptors; +} prof_hsa_loader_pfn_t; -static const char *getDeviceArchName(int DeviceId) { - if (DeviceId < 0 || DeviceId >= NumDevices || !DeviceArchNames[DeviceId][0]) - return "amdgpu"; - return DeviceArchNames[DeviceId]; -} +/* HIP: only hipMemcpy is needed, for device-to-host copies. */ +typedef int (*hipMemcpy_ty)(void *, const void *, size_t, int); /* -------------------------------------------------------------------------- */ -/* Dynamic module tracking */ +/* Resolved runtime entry points */ /* -------------------------------------------------------------------------- */ -/* Per-TU profile entry inside a dynamic module. - * A single dynamic module may contain multiple TUs (e.g. -fgpu-rdc). */ -typedef struct { - void *DeviceVar; /* device address of __llvm_profile_sections_ */ - int Processed; /* 0 = not yet collected, 1 = data already copied */ -} OffloadDynamicTUInfo; +static hsa_iterate_agents_ty pHsaIterateAgents = nullptr; +static hsa_agent_get_info_ty pHsaAgentGetInfo = nullptr; +static hsa_executable_iterate_agent_symbols_ty pHsaExecIterAgentSyms = nullptr; +static hsa_executable_symbol_get_info_ty pHsaSymGetInfo = nullptr; +static hsa_loader_query_segment_descriptors_ty pQuerySegDescs = nullptr; +static hipMemcpy_ty pHipMemcpy = nullptr; -/* One entry per hipModuleLoad call. */ -typedef struct { - void *ModulePtr; /* hipModule_t handle */ - OffloadDynamicTUInfo *TUs; /* array of per-TU entries */ - int NumTUs; - int CapTUs; -} OffloadDynamicModuleInfo; +/* 0 = not yet attempted, 1 = ready, -1 = unavailable. + * Accessed with acquire/release atomics: a thread that observes RuntimeState==1 + * (acquire) is guaranteed to also see the fully-written p* function pointers + * (which are published before the release store of RuntimeState=1 below). */ +static int RuntimeState = 0; -static OffloadDynamicModuleInfo *DynamicModules = nullptr; -static int NumDynamicModules = 0; -static int CapDynamicModules = 0; +static int isVerboseMode(void) { + static int IsVerbose = -1; + if (IsVerbose == -1) + IsVerbose = getenv("LLVM_PROFILE_VERBOSE") != nullptr; + return IsVerbose; +} /* -------------------------------------------------------------------------- */ -/* ELF symbol enumeration (manual parse: compiler-rt cannot link LLVM Support) - */ +/* One-time runtime resolution */ /* -------------------------------------------------------------------------- */ -#if __has_include() -#include - -/* Callback invoked for every matching symbol name found in the ELF image. - * Return 0 to continue iteration, non-zero to stop. */ -typedef int (*SymbolCallback)(const char *Name, void *UserData); - -/* If Image is a clang offload bundle, return a pointer to the first embedded - * ELF. Returns Image if not a bundle, nullptr if a bundle holds no ELF. */ -static const void *unwrapOffloadBundle(const void *Image) { - static const char BundleMagic[] = "__CLANG_OFFLOAD_BUNDLE__"; - if (memcmp(Image, BundleMagic, sizeof(BundleMagic) - 1) != 0) - return Image; /* Not a bundle, return as-is. */ - - const char *Buf = (const char *)Image; - uint64_t NumEntries; - __builtin_memcpy(&NumEntries, Buf + sizeof(BundleMagic) - 1, - sizeof(uint64_t)); - - /* Walk the entry table (starts at offset 32). */ - const char *Cursor = Buf + 32; - for (uint64_t I = 0; I < NumEntries; ++I) { - uint64_t EntryOffset, EntrySize, IDSize; - __builtin_memcpy(&EntryOffset, Cursor, sizeof(EntryOffset)); - Cursor += sizeof(EntryOffset); - __builtin_memcpy(&EntrySize, Cursor, sizeof(EntrySize)); - Cursor += sizeof(EntrySize); - __builtin_memcpy(&IDSize, Cursor, sizeof(IDSize)); - Cursor += sizeof(IDSize); - Cursor += IDSize; /* skip entry ID */ - - if (EntrySize >= sizeof(Elf64_Ehdr)) { - const Elf64_Ehdr *E = (const Elf64_Ehdr *)(Buf + EntryOffset); - if (E->e_ident[EI_MAG0] == ELFMAG0 && E->e_ident[EI_MAG1] == ELFMAG1 && - E->e_ident[EI_MAG2] == ELFMAG2 && E->e_ident[EI_MAG3] == ELFMAG3) { - return (const void *)(Buf + EntryOffset); - } - } - } - - PROF_WARN("%s", "offload bundle contains no valid ELF entries\n"); - return nullptr; +/* Publish the terminal RuntimeState (release) and map it to the 0/-1 return + * convention used by loadRuntimePointers(). */ +static int setRuntimeState(int S) { + __atomic_store_n(&RuntimeState, S, __ATOMIC_RELEASE); + return S > 0 ? 0 : -1; } -/* Invoke CB for every global symbol in Image (an AMDGPU ELF or offload bundle) - * whose name starts with PREFIX. Image may be null. */ -static void enumerateElfSymbols(const void *Image, const char *Prefix, - SymbolCallback CB, void *UserData) { - if (!Image) - return; - - Image = unwrapOffloadBundle(Image); - if (!Image) - return; +static int loadRuntimePointers(void) { + int State = __atomic_load_n(&RuntimeState, __ATOMIC_ACQUIRE); + if (State) + return State > 0 ? 0 : -1; - 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 (!__interception::DynamicLoaderAvailable()) { if (isVerboseMode()) - PROF_NOTE("%s", "Image is not a valid ELF, skipping enumeration\n"); - return; + PROF_NOTE("%s", "Dynamic library loading not available - " + "device profiling disabled\n"); + return setRuntimeState(-1); } - 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; - } - } + void *Hsa = __interception::OpenLibrary("libhsa-runtime64.so"); + if (!Hsa) + Hsa = __interception::OpenLibrary("libhsa-runtime64.so.1"); + if (!Hsa) { + if (isVerboseMode()) + PROF_NOTE("%s", "libhsa-runtime64.so not loadable - " + "device profiling disabled\n"); + return setRuntimeState(-1); } -} -/* State passed through the enumeration callback. */ -typedef struct { - void *Module; /* hipModule_t */ - OffloadDynamicModuleInfo *ModInfo; -} EnumState; - -/* Register one __llvm_profile_sections_ symbol on the module entry. - * hipModuleGetGlobal also registers the device address with CLR so hipMemcpy - * can copy from it later. */ -static int registerPrfSymbol(const char *Name, void *UserData) { - EnumState *S = (EnumState *)UserData; - OffloadDynamicModuleInfo *MI = S->ModInfo; - - /* The symbol is the per-TU sections struct itself, not a pointer - * indirection, so this address is the hipMemcpy source. */ - void *DeviceVar = nullptr; - size_t Bytes = 0; - if (hipModuleGetGlobal(&DeviceVar, &Bytes, S->Module, Name) != 0) { - PROF_WARN("failed to get symbol %s for module %p\n", Name, S->Module); - return 0; /* continue */ - } - - if (MI->NumTUs >= MI->CapTUs) { - int NewCap = MI->CapTUs ? MI->CapTUs * 2 : 4; - OffloadDynamicTUInfo *New = (OffloadDynamicTUInfo *)realloc( - MI->TUs, NewCap * sizeof(OffloadDynamicTUInfo)); - if (!New) { - PROF_ERR("%s\n", "failed to grow TU array"); - return 0; - } - MI->TUs = New; - MI->CapTUs = NewCap; + hsa_init_ty pHsaInit = + (hsa_init_ty)__interception::LookupSymbol(Hsa, "hsa_init"); + hsa_system_get_major_extension_table_ty pGetExtTable = + (hsa_system_get_major_extension_table_ty)__interception::LookupSymbol( + Hsa, "hsa_system_get_major_extension_table"); + pHsaIterateAgents = (hsa_iterate_agents_ty)__interception::LookupSymbol( + Hsa, "hsa_iterate_agents"); + pHsaAgentGetInfo = (hsa_agent_get_info_ty)__interception::LookupSymbol( + Hsa, "hsa_agent_get_info"); + pHsaExecIterAgentSyms = + (hsa_executable_iterate_agent_symbols_ty)__interception::LookupSymbol( + Hsa, "hsa_executable_iterate_agent_symbols"); + pHsaSymGetInfo = + (hsa_executable_symbol_get_info_ty)__interception::LookupSymbol( + Hsa, "hsa_executable_symbol_get_info"); + + if (!pHsaInit || !pGetExtTable || !pHsaIterateAgents || !pHsaAgentGetInfo || + !pHsaExecIterAgentSyms || !pHsaSymGetInfo) { + PROF_WARN("%s", "required HSA symbols missing - device profiling disabled\n"); + return setRuntimeState(-1); } - OffloadDynamicTUInfo *TU = &MI->TUs[MI->NumTUs++]; - TU->DeviceVar = DeviceVar; - TU->Processed = 0; - - (void)Name; - return 0; /* continue enumeration */ -} -#endif /* __has_include() */ - -/* -------------------------------------------------------------------------- */ -/* Registration / un-registration helpers */ -/* -------------------------------------------------------------------------- */ - -extern "C" void -__llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr, - const void *Image) { - if (ModuleLoadRc) - return; - - lockDynamicModules(); - - if (isVerboseMode()) - PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n", - NumDynamicModules, ModuleLoadRc, *Ptr, Image); - - if (NumDynamicModules >= CapDynamicModules) { - int NewCap = CapDynamicModules ? CapDynamicModules * 2 : 64; - OffloadDynamicModuleInfo *New = (OffloadDynamicModuleInfo *)realloc( - DynamicModules, NewCap * sizeof(OffloadDynamicModuleInfo)); - if (!New) { - unlockDynamicModules(); - return; - } - DynamicModules = New; - CapDynamicModules = NewCap; + /* Bring HSA up now (idempotent, refcounted). Doing this in the library + * constructor guarantees HSA registers its own atexit(hsa_shut_down) + * before we register atexit(drainDevices); atexit is LIFO, so our drain + * runs while HSA is still alive. */ + prof_hsa_status_t St = pHsaInit(); + if (St != PROF_HSA_STATUS_SUCCESS && St != PROF_HSA_STATUS_INFO_BREAK) { + if (isVerboseMode()) + PROF_NOTE("hsa_init failed (0x%x) - device profiling disabled\n", St); + return setRuntimeState(-1); } - OffloadDynamicModuleInfo *MI = &DynamicModules[NumDynamicModules++]; - MI->ModulePtr = *Ptr; - MI->TUs = nullptr; - MI->NumTUs = 0; - MI->CapTUs = 0; - - /* Dynamic-module profiling needs ELF parsing for symbol enumeration. */ -#if __has_include() - EnumState State = {*Ptr, MI}; - enumerateElfSymbols(Image, "__llvm_profile_sections_", registerPrfSymbol, - &State); + prof_hsa_loader_pfn_t LoaderApi; + __builtin_memset(&LoaderApi, 0, sizeof(LoaderApi)); + St = pGetExtTable(PROF_HSA_EXTENSION_AMD_LOADER, 1, sizeof(LoaderApi), + &LoaderApi); + if (St != PROF_HSA_STATUS_SUCCESS || !LoaderApi.query_segment_descriptors) { + PROF_WARN("AMD loader extension unavailable (0x%x) - " + "device profiling disabled\n", + St); + return setRuntimeState(-1); + } + pQuerySegDescs = LoaderApi.query_segment_descriptors; + + /* HIP lookup is best-effort across deployment shapes: + * 1. The vast majority of HIP-using programs already have libamdhip64 + * loaded (the application or one of its libraries linked it directly). + * Resolving via the process namespace catches that case without us + * having to know the SONAME, and works even when there is no + * development "libamdhip64.so" symlink (runtime-only ROCm installs). + * 2. If hipMemcpy is not in the namespace, fall back to dlopen, trying + * versioned SONAMEs first (which is what the dynamic linker actually + * loads at program start) before the unversioned dev symlink. + * 3. On Windows just dlopen the DLL by name. */ + pHipMemcpy = + (hipMemcpy_ty)__interception::LookupSymbolDefault("hipMemcpy"); + if (!pHipMemcpy) { +#ifdef _WIN32 + static const char *const HipLibNames[] = {"amdhip64.dll", nullptr}; #else - (void)Image; - if (isVerboseMode()) - PROF_NOTE("%s", - "Dynamic module profiling not supported on this platform\n"); + /* Order: most recent ROCm major first, then older majors, then the + * unversioned development symlink as a last resort. Update this list + * when ROCm bumps libamdhip64 SONAME. */ + static const char *const HipLibNames[] = { + "libamdhip64.so.7", "libamdhip64.so.6", "libamdhip64.so.5", + "libamdhip64.so.4", "libamdhip64.so", nullptr}; #endif - - if (MI->NumTUs == 0) { - PROF_WARN("no __llvm_profile_sections_* symbols found in module %p\n", - *Ptr); - } else if (isVerboseMode()) { - PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs); - } - - unlockDynamicModules(); -} - -extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) { - lockDynamicModules(); - for (int i = 0; i < NumDynamicModules; ++i) { - OffloadDynamicModuleInfo *MI = &DynamicModules[i]; - - /* HIP recycles hipModule_t addresses; drained slots are cleared so a - * recycled handle finds the new slot, not the dead one. */ - if (MI->ModulePtr != Ptr) - continue; - - if (isVerboseMode()) - PROF_NOTE("Unregistering module %p (%d TUs)\n", MI->ModulePtr, - MI->NumTUs); - - static int NextTUIndex = 0; - 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); + for (int i = 0; HipLibNames[i] != nullptr; ++i) { + void *Hip = __interception::OpenLibrary(HipLibNames[i]); + if (!Hip) continue; - } - int TUIndex = __atomic_fetch_add(&NextTUIndex, 1, __ATOMIC_RELAXED); - if (TU->DeviceVar) { - int CurDev = 0; - hipGetDevice(&CurDev); - const char *ArchName = getDeviceArchName(CurDev); - /* Encode TUIndex in Target so each drain writes a distinct profraw; - * otherwise back-to-back drains overwrite the same file. */ - char TargetWithTU[64]; - snprintf(TargetWithTU, sizeof(TargetWithTU), "%s.%d", ArchName, - TUIndex); - if (processDeviceOffloadPrf(TU->DeviceVar, TUIndex, TargetWithTU) == 0) - TU->Processed = 1; - else - PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr, - t); + pHipMemcpy = + (hipMemcpy_ty)__interception::LookupSymbol(Hip, "hipMemcpy"); + if (pHipMemcpy) { + if (isVerboseMode()) + PROF_NOTE("HIP resolved via dlopen(%s)\n", HipLibNames[i]); + break; } } - MI->ModulePtr = nullptr; - unlockDynamicModules(); - return; + } else if (isVerboseMode()) { + PROF_NOTE("%s", + "HIP resolved via existing process namespace (RTLD_DEFAULT)\n"); + } + if (!pHipMemcpy) { + PROF_WARN("%s", "hipMemcpy unavailable - device profiling disabled\n"); + return setRuntimeState(-1); } if (isVerboseMode()) - PROF_WARN("unregister called for unknown module %p\n", Ptr); - unlockDynamicModules(); -} - -/* 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; + PROF_NOTE("%s", "HSA + HIP runtime resolved for device profiling\n"); + return setRuntimeState(1); } -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 int memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) { + return pHipMemcpy ? pHipMemcpy(Dst, Src, Size, 2 /* hipMemcpyDeviceToHost */) + : -1; } -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()-based scope guard */ +/* -------------------------------------------------------------------------- */ namespace { - -// free()-based scope guard. Use .release() to transfer ownership. struct UniqueFree { void *Ptr; explicit UniqueFree(void *P = nullptr) : Ptr(P) {} @@ -538,35 +311,78 @@ struct UniqueFree { free(Ptr); Ptr = P; } - void *release() { - void *P = Ptr; - Ptr = nullptr; - return P; - } }; - } // namespace -static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex, - const char *Target) { - __llvm_profile_gpu_sections HostSections; +/* -------------------------------------------------------------------------- */ +/* Copy one device bounds table to the host and emit a .profraw */ +/* -------------------------------------------------------------------------- */ + +/* Plausibility cap for any single device-side profile section. Device + * profile data for a single linked code object is typically <10 MiB; a few + * hundred MiB would already be unprecedented. Setting this to 256 MiB + * catches a corrupted/uninitialized bounds table early (where End-Begin + * can compute to multi-GiB) before we try to malloc/memcpy bogus memory. */ +#define PROF_MAX_SECTION_BYTES ((size_t)256 * 1024 * 1024) + +/* uintptr_t-based size of a [Begin, End] device range, with validation. + * Returns -1 if End < Begin (would wrap to a huge size_t) or if the span + * exceeds the per-section cap. On success, *OutSize is the byte count. */ +static int computeRangeSize(const char *Label, const void *Begin, + const void *End, size_t *OutSize) { + uintptr_t B = (uintptr_t)Begin; + uintptr_t E = (uintptr_t)End; + if (E < B) { + PROF_WARN("device %s range invalid: end %p < begin %p\n", Label, End, + Begin); + return -1; + } + size_t Sz = (size_t)(E - B); + if (Sz > PROF_MAX_SECTION_BYTES) { + PROF_WARN("device %s range %zu bytes exceeds %zu-byte cap; refusing " + "to copy (likely corrupted bounds table)\n", + Label, Sz, (size_t)PROF_MAX_SECTION_BYTES); + return -1; + } + *OutSize = Sz; + return 0; +} - if (hipMemcpy(&HostSections, DeviceOffloadPrf, sizeof(HostSections), - 2 /*DToH*/) != 0) { - PROF_ERR("%s\n", "failed to copy offload prf structure from device"); +/* Returns 1 if a device .profraw was written, 0 if there was nothing to write + * (empty counters/data sections), and -1 on error. The caller distinguishes + * these so an empty section is never miscounted as a successful drain. */ +static int processDeviceSections(void *DeviceSectionsAddr, const char *Target) { + __llvm_profile_gpu_sections HostSections; + if (memcpyDeviceToHost(&HostSections, DeviceSectionsAddr, + sizeof(HostSections)) != 0) { + PROF_ERR("%s\n", "failed to copy device bounds table 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 *DevDataBegin = HostSections.DataStart; const void *DevDataEnd = HostSections.DataStop; + const void *DevNamesBegin = HostSections.NamesStart; 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; + size_t CountersSize, DataSize, NamesSize; + if (computeRangeSize("counters", DevCntsBegin, DevCntsEnd, &CountersSize) != + 0 || + computeRangeSize("data", DevDataBegin, DevDataEnd, &DataSize) != 0 || + computeRangeSize("names", DevNamesBegin, DevNamesEnd, &NamesSize) != 0) + return -1; + + /* DataSize must be an integral number of __llvm_profile_data records; + * otherwise either the table layout has changed under us or the bounds + * point at the wrong section. Refuse to relocate it - the per-record + * loop below would walk off the end. */ + if (DataSize % sizeof(__llvm_profile_data) != 0) { + PROF_WARN("device data section size %zu is not a multiple of " + "sizeof(__llvm_profile_data)=%zu\n", + DataSize, sizeof(__llvm_profile_data)); + return -1; + } if (isVerboseMode()) PROF_NOTE("Section pointers: Cnts=[%p,%p]=%zu Data=[%p,%p]=%zu " @@ -577,164 +393,110 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex, 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; - - // Owns freshly malloc'd buffers; release() transfers ownership to the cache. UniqueFree CntsOwner, DataOwner, NamesOwner; - - 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); - CntsOwner.reset(HostCountersBegin); - } - - 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); - DataOwner.reset(HostDataBegin); - } - - 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); - NamesOwner.reset(HostNamesBegin); - } - - if ((DataSize > 0 && !HostDataBegin) || - (CountersSize > 0 && !HostCountersBegin) || - (NamesSize > 0 && !HostNamesBegin)) { + char *HostCounters = (char *)malloc(CountersSize); + CntsOwner.reset(HostCounters); + char *HostData = (char *)malloc(DataSize); + DataOwner.reset(HostData); + char *HostNames = NamesSize ? (char *)malloc(NamesSize) : nullptr; + if (NamesSize) + NamesOwner.reset(HostNames); + + if (!HostCounters || !HostData || (NamesSize && !HostNames)) { PROF_ERR("%s\n", "failed to allocate host memory for device sections"); return -1; } - 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)) { + if (memcpyDeviceToHost(HostData, DevDataBegin, DataSize) != 0 || + memcpyDeviceToHost(HostCounters, DevCntsBegin, CountersSize) != 0 || + (NamesSize && + memcpyDeviceToHost(HostNames, DevNamesBegin, NamesSize) != 0)) { PROF_ERR("%s\n", "failed to copy profile sections from device"); return -1; } - /* Cache buffers so RDC-mode multi-shadow drains can reuse them. - * release() prevents the scope guards from freeing what the cache owns. */ - if (!CntsReused && CountersSize > 0) { - CachedDevCntsBegin = DevCntsBegin; - CachedHostCnts = HostCountersBegin; - CachedCntsSize = CountersSize; - CntsOwner.release(); - } - if (!DataReused && DataSize > 0) { - CachedDevDataBegin = DevDataBegin; - CachedHostData = HostDataBegin; - CachedDataSize = DataSize; - DataOwner.release(); - } - if (!NamesReused && NamesSize > 0) { - CachedDevNamesBegin = DevNamesBegin; - CachedHostNames = HostNamesBegin; - CachedNamesSize = NamesSize; - NamesOwner.release(); - } - 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). + // Lay the buffer out as [Data][PaddingBeforeCounters][Counters][Names] to + // match what lprofWriteDataImpl expects (CountersDelta = Counters - Data). 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; - + uint64_t PadBeforeCounters, PadAfterCounters, PadAfterBitmap, PadAfterNames, + PadAfterVTable, PadAfterVNames; if (__llvm_profile_get_padding_sizes_for_counters( - DataSize, CountersSize, NumBitmapBytes, NamesSize, VTableSectionSize, - VNamesSize, &PaddingBytesBeforeCounters, &PaddingBytesAfterCounters, - &PaddingBytesAfterBitmapBytes, &PaddingBytesAfterNames, - &PaddingBytesAfterVTable, &PaddingBytesAfterVNames) != 0) { + DataSize, CountersSize, /*NumBitmapBytes=*/0, NamesSize, + /*VTableSize=*/0, /*VNameSize=*/0, &PadBeforeCounters, + &PadAfterCounters, &PadAfterBitmap, &PadAfterNames, &PadAfterVTable, + &PadAfterVNames) != 0) { PROF_ERR("%s\n", "failed to get padding sizes"); return -1; } - size_t ContiguousBufferSize = - DataSize + PaddingBytesBeforeCounters + CountersSize + NamesSize; - UniqueFree ContiguousBuf(malloc(ContiguousBufferSize)); - if (!ContiguousBuf.get()) { + size_t BufSize = DataSize + PadBeforeCounters + CountersSize + NamesSize; + UniqueFree BufOwner(malloc(BufSize)); + char *Buf = BufOwner.get(); + if (!Buf) { PROF_ERR("%s\n", "failed to allocate contiguous buffer"); return -1; } - char *ContiguousBuffer = ContiguousBuf.get(); - __builtin_memset(ContiguousBuffer, 0, ContiguousBufferSize); - - char *BufDataBegin = ContiguousBuffer; - char *BufCountersBegin = - ContiguousBuffer + DataSize + PaddingBytesBeforeCounters; - char *BufNamesBegin = BufCountersBegin + CountersSize; - - __builtin_memcpy(BufDataBegin, HostDataBegin, DataSize); - __builtin_memcpy(BufCountersBegin, HostCountersBegin, CountersSize); - __builtin_memcpy(BufNamesBegin, HostNamesBegin, NamesSize); - - // CounterPtr is a device-relative offset; relocate it for the file layout - // where the Data section precedes Counters. - __llvm_profile_data *RelocatedData = (__llvm_profile_data *)BufDataBegin; + __builtin_memset(Buf, 0, BufSize); + + char *BufData = Buf; + char *BufCounters = Buf + DataSize + PadBeforeCounters; + char *BufNames = BufCounters + CountersSize; + + __builtin_memcpy(BufData, HostData, DataSize); + __builtin_memcpy(BufCounters, HostCounters, CountersSize); + if (NamesSize) + __builtin_memcpy(BufNames, HostNames, NamesSize); + + // Relocate each record's CounterPtr from the device-relative offset to the + // file-layout-relative offset (Data section precedes Counters in the file). + // Validate every resolved device counter address lies within the copied + // counters region; out-of-range entries indicate a stale/mismatched bounds + // table and would otherwise produce a .profraw with counters pointing at + // unrelated memory. + __llvm_profile_data *RelocatedData = (__llvm_profile_data *)BufData; + int BadRecords = 0; 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)); - __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) + - offsetof(__llvm_profile_data, CounterPtr), - &NewRelativeOffset, sizeof(NewRelativeOffset)); + uintptr_t DeviceDataStructAddr = + (uintptr_t)DevDataBegin + (uintptr_t)(i * sizeof(__llvm_profile_data)); + uintptr_t DeviceCountersAddr = + DeviceDataStructAddr + (uintptr_t)DeviceCounterPtrOffset; + uintptr_t CntsB = (uintptr_t)DevCntsBegin; + uintptr_t CntsE = (uintptr_t)DevCntsEnd; + /* Allow CountersAddr == CntsE for a zero-counter record at the very + * end of the section. */ + if (DeviceCountersAddr < CntsB || DeviceCountersAddr > CntsE) { + BadRecords++; + if (isVerboseMode()) + PROF_NOTE("record %llu: device counter addr %p outside " + "[%p,%p]; zeroing CounterPtr\n", + (unsigned long long)i, (void *)DeviceCountersAddr, + DevCntsBegin, DevCntsEnd); + // CounterPtr is IntPtrT (pointer-sized): zero exactly that field so we + // never clobber adjacent record fields on a 32-bit host. + __builtin_memset((char *)RelocatedData + + i * sizeof(__llvm_profile_data) + + offsetof(__llvm_profile_data, CounterPtr), + 0, sizeof(RelocatedData[i].CounterPtr)); + } else { + ptrdiff_t OffsetIntoCountersSection = + (ptrdiff_t)(DeviceCountersAddr - CntsB); + ptrdiff_t NewRelativeOffset = + (ptrdiff_t)DataSize + (ptrdiff_t)PadBeforeCounters + + OffsetIntoCountersSection - + (ptrdiff_t)(i * sizeof(__llvm_profile_data)); + __builtin_memcpy((char *)RelocatedData + + i * sizeof(__llvm_profile_data) + + offsetof(__llvm_profile_data, CounterPtr), + &NewRelativeOffset, sizeof(NewRelativeOffset)); + } } + // Zero the fields the writer does not expect to be populated. __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) + offsetof(__llvm_profile_data, BitmapPtr), 0, @@ -742,156 +504,395 @@ static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex, sizeof(RelocatedData[i].FunctionPointer) + sizeof(RelocatedData[i].Values)); } + if (BadRecords > 0) + PROF_WARN("%d/%llu device profile record(s) had out-of-range " + "counter pointers (zeroed)\n", + BadRecords, (unsigned long long)NumData); + + int Ret = __llvm_write_custom_profile( + Target, (__llvm_profile_data *)BufData, + (__llvm_profile_data *)(BufData + DataSize), BufCounters, + BufCounters + CountersSize, BufNames, BufNames + NamesSize, nullptr); + + if (Ret != 0) { + PROF_ERR("%s\n", "failed to write device profile"); + return -1; + } + if (isVerboseMode()) + PROF_NOTE("Wrote device profile (target=%s)\n", Target); + return 1; +} - /* Target already encodes TUIndex when needed. */ - (void)TUIndex; +/* -------------------------------------------------------------------------- */ +/* HSA walk */ +/* -------------------------------------------------------------------------- */ - ret = __llvm_write_custom_profile( - Target, (__llvm_profile_data *)BufDataBegin, - (__llvm_profile_data *)(BufDataBegin + DataSize), BufCountersBegin, - BufCountersBegin + CountersSize, BufNamesBegin, BufNamesBegin + NamesSize, - nullptr); +#define PROF_MAX_GPU_AGENTS 64 - 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"); - } +namespace { +struct GpuAgent { + prof_hsa_agent_t agent; + char arch[64]; +}; + +struct WalkState { + GpuAgent agents[PROF_MAX_GPU_AGENTS]; + int num_agents; + int total_found; + int total_drained; +}; + +/* Per (agent, executable) symbol-iteration state. */ +struct SymbolState { + const char *arch; + int found; + int drained; +}; +} // namespace + +/* The canonical device bounds table symbol from InstrProfilingPlatformGPU.c. */ +static const char ProfileSectionsSymbol[] = "__llvm_profile_sections"; + +/* Dedup distinct (Data,Counters,Names) tuples: a single linked device code + * object exposes one __llvm_profile_sections, but the same bounds may be seen + * via multiple agents, so drain each unique counter set only once. Also used + * to generate collision-free target names. */ +namespace { +struct BoundsTuple { + const void *data; + const void *cnts; + const void *names; +}; +} // namespace - return ret; +#define PROF_MAX_SEEN_BOUNDS 256 +static BoundsTuple SeenBounds[PROF_MAX_SEEN_BOUNDS]; +static int NumSeenBounds = 0; + +/* Pure check: has this bounds tuple already been processed? Does not mutate + * state, so a transient drain failure does not permanently suppress retries. */ +static int seenBounds(const void *D, const void *C, const void *N) { + for (int i = 0; i < NumSeenBounds; ++i) + if (SeenBounds[i].data == D && SeenBounds[i].cnts == C && + SeenBounds[i].names == N) + return 1; + return 0; } -static int processShadowVariable(void *ShadowVar, int TUIndex, - const char *Target) { - void *DeviceSections = nullptr; - if (hipGetSymbolAddress(&DeviceSections, ShadowVar) != 0) { - PROF_WARN("failed to get symbol address for shadow variable %p\n", - ShadowVar); - return -1; +/* Record a bounds tuple as processed. Only called after a successful drain or a + * confirmed-empty section, so failed attempts stay retryable. Idempotent. */ +static void recordBounds(const void *D, const void *C, const void *N) { + if (seenBounds(D, C, N)) + return; + if (NumSeenBounds < PROF_MAX_SEEN_BOUNDS) { + SeenBounds[NumSeenBounds].data = D; + SeenBounds[NumSeenBounds].cnts = C; + SeenBounds[NumSeenBounds].names = N; + NumSeenBounds++; } - /* DeviceSections points at the per-TU sections struct itself. */ - return processDeviceOffloadPrf(DeviceSections, TUIndex, Target); } -static int isHipAvailable(void) { - ensureHipLoaded(); - return pHipMemcpy != nullptr && pHipGetSymbolAddress != nullptr; +static prof_hsa_status_t onSymbol(prof_hsa_executable_t, prof_hsa_agent_t, + prof_hsa_executable_symbol_t Sym, + void *Data) { + SymbolState *S = (SymbolState *)Data; + + prof_hsa_symbol_kind_t Kind; + if (pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &Kind) != + PROF_HSA_STATUS_SUCCESS || + Kind != PROF_HSA_SYMBOL_KIND_VARIABLE) + return PROF_HSA_STATUS_SUCCESS; + + uint32_t NameLen = 0; + if (pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, + &NameLen) != PROF_HSA_STATUS_SUCCESS || + NameLen != sizeof(ProfileSectionsSymbol) - 1) + return PROF_HSA_STATUS_SUCCESS; + + char NameBuf[64]; + if (NameLen + 1 > sizeof(NameBuf)) + return PROF_HSA_STATUS_SUCCESS; + if (pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_NAME, NameBuf) != + PROF_HSA_STATUS_SUCCESS) + return PROF_HSA_STATUS_SUCCESS; + NameBuf[NameLen] = '\0'; + + if (strcmp(NameBuf, ProfileSectionsSymbol) != 0) + return PROF_HSA_STATUS_SUCCESS; + + uint64_t Addr = 0; + if (pHsaSymGetInfo(Sym, PROF_HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, + &Addr) != PROF_HSA_STATUS_SUCCESS || + Addr == 0) { + if (isVerboseMode()) + PROF_NOTE("%s", "failed to read __llvm_profile_sections address\n"); + return PROF_HSA_STATUS_SUCCESS; + } + + S->found++; + + // Read the bounds table first to dedup before the full copy. + __llvm_profile_gpu_sections Sec; + if (memcpyDeviceToHost(&Sec, (void *)(uintptr_t)Addr, sizeof(Sec)) != 0) { + PROF_WARN("%s", "failed to copy device bounds table\n"); + return PROF_HSA_STATUS_SUCCESS; + } + if (seenBounds(Sec.DataStart, Sec.CountersStart, Sec.NamesStart)) { + if (isVerboseMode()) + PROF_NOTE("%s", "device bounds already drained, skipping\n"); + return PROF_HSA_STATUS_SUCCESS; + } + + // Generate a collision-free target. Multiple distinct device code objects on + // the same arch (e.g. non-RDC multi-TU) must not clobber each other's file. + static int DrainIndex = 0; + char Target[96]; + if (DrainIndex == 0) + snprintf(Target, sizeof(Target), "%s", S->arch); + else + snprintf(Target, sizeof(Target), "%s.%d", S->arch, DrainIndex); + + // Only a >0 result means a .profraw was actually written; an empty section + // (0) or an error (<0) must not be counted as a drain or advance DrainIndex. + // Record the bounds as processed only on success (>0) or a confirmed-empty + // section (0); a transient error (<0) is left unrecorded so a later agent or + // a subsequent collect call can retry instead of silently dropping data. + int Rc = processDeviceSections((void *)(uintptr_t)Addr, Target); + if (Rc > 0) { + S->drained++; + DrainIndex++; + recordBounds(Sec.DataStart, Sec.CountersStart, Sec.NamesStart); + } else if (Rc == 0) { + recordBounds(Sec.DataStart, Sec.CountersStart, Sec.NamesStart); + } + + return PROF_HSA_STATUS_SUCCESS; } -/* -------------------------------------------------------------------------- */ -/* Collect device-side profile data */ -/* -------------------------------------------------------------------------- */ +static prof_hsa_status_t collectAgent(prof_hsa_agent_t Agent, void *Data) { + prof_hsa_device_type_t DevType; + if (pHsaAgentGetInfo(Agent, PROF_HSA_AGENT_INFO_DEVICE, &DevType) != + PROF_HSA_STATUS_SUCCESS || + DevType != PROF_HSA_DEVICE_TYPE_GPU) + return PROF_HSA_STATUS_SUCCESS; -extern "C" int __llvm_profile_hip_collect_device_data(void) { - if (NumShadowVariables == 0 && NumDynamicModules == 0) + WalkState *W = (WalkState *)Data; + if (W->num_agents >= PROF_MAX_GPU_AGENTS) + return PROF_HSA_STATUS_SUCCESS; + + GpuAgent &GA = W->agents[W->num_agents++]; + GA.agent = Agent; + char Name[64]; + __builtin_memset(Name, 0, sizeof(Name)); + pHsaAgentGetInfo(Agent, PROF_HSA_AGENT_INFO_NAME, Name); + size_t N = strnlen(Name, sizeof(GA.arch) - 1); + __builtin_memcpy(GA.arch, Name, N); + GA.arch[N] = '\0'; + if (!GA.arch[0]) + strncpy(GA.arch, "amdgpu", sizeof(GA.arch) - 1); + + if (isVerboseMode()) + PROF_NOTE("GPU agent %d: %s\n", W->num_agents - 1, GA.arch); + return PROF_HSA_STATUS_SUCCESS; +} + +/* Reentrancy guard and "we drained data at least once" flag. Both the host + * write path and the atexit handler call drainDevices(); a successful walk + * with non-empty results latches DrainCompleted so we never re-emit duplicate + * .profraw files, but transient no-op outcomes ("runtime not yet loadable", + * "no GPU agents", "no loaded segments", "no instrumented sections found") + * stay retryable so the final atexit drain can still pick up code objects + * that loaded later. The InProgress flag prevents a concurrent call from + * another thread (or a re-entrant call on the same thread, e.g. a library + * destructor that triggers another drain) from running the walk concurrently + * and corrupting the global SeenBounds table. Both flags are accessed with + * acquire/release atomics so the guard holds across threads. */ +static int DrainInProgress = 0; +static int DrainCompleted = 0; + +static int drainDevices(void) { + if (__atomic_load_n(&DrainCompleted, __ATOMIC_ACQUIRE)) return 0; - if (!isHipAvailable()) + /* Claim the drain with an atomic CAS. A failed CAS means either another + * thread is already draining, or this is a reentrant call on the same + * thread (e.g. a library destructor that triggers another drain); both + * must bail without touching the global SeenBounds table. The acquire/ + * release ordering also publishes the worker thread's writes to threads + * that observe DrainCompleted later. */ + int Expected = 0; + if (!__atomic_compare_exchange_n(&DrainInProgress, &Expected, 1, + /*weak=*/0, __ATOMIC_ACQ_REL, + __ATOMIC_ACQUIRE)) return 0; - int Ret = 0; + /* Mirror the early-exit paths so we always release the in-progress flag. */ + struct InProgressGuard { + ~InProgressGuard() { + __atomic_store_n(&DrainInProgress, 0, __ATOMIC_RELEASE); + } + } _Guard; - /* Shadow variables (static-linked kernels): drain from every device. */ - if (NumShadowVariables > 0) { - int OrigDevice = -1; - hipGetDevice(&OrigDevice); + if (loadRuntimePointers() != 0) { + /* Runtime unavailable: don't latch DrainCompleted, allow a later call + * (e.g. atexit, after the host has dlopen'd HIP) to retry. */ + return 0; + } - 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) { - /* RDC-mode multi-shadow drains need a distinct profraw per TU; - * single-TU programs keep the bare arch target. */ - const char *Target = ArchName; - char TargetWithIdx[64]; - if (NumShadowVariables > 1) { - snprintf(TargetWithIdx, sizeof(TargetWithIdx), "%s.%d", ArchName, i); - Target = TargetWithIdx; - } - if (processShadowVariable(OffloadShadowVariables[i], i, Target) != 0) - Ret = -1; - } - } + WalkState W; + __builtin_memset(&W, 0, sizeof(W)); + prof_hsa_status_t St = pHsaIterateAgents(collectAgent, &W); + if (St != PROF_HSA_STATUS_SUCCESS && St != PROF_HSA_STATUS_INFO_BREAK) { + PROF_WARN("hsa_iterate_agents failed (0x%x)\n", St); + return -1; + } + if (W.num_agents == 0) { + if (isVerboseMode()) + PROF_NOTE("%s", "no GPU agents present; nothing to drain (will retry)\n"); + return 0; + } - if (OrigDevice >= 0) - hipSetDevice(OrigDevice); + /* query_segment_descriptors ships in every loader-extension version and is + * more permissive than iterate_executables on ROCm. It yields the loaded + * (agent, executable) pairs directly. */ + size_t NumSegs = 0; + St = pQuerySegDescs(nullptr, &NumSegs); + if (St != PROF_HSA_STATUS_SUCCESS) { + PROF_WARN("query_segment_descriptors(count) failed (0x%x)\n", St); + return -1; + } + if (NumSegs == 0) { + if (isVerboseMode()) + PROF_NOTE("%s", + "no loaded segments; nothing to drain (will retry)\n"); + return 0; + } + + prof_hsa_loader_segment_descriptor_t *Segs = + (prof_hsa_loader_segment_descriptor_t *)calloc(NumSegs, sizeof(*Segs)); + if (!Segs) { + PROF_ERR("%s\n", "failed to allocate segment descriptor array"); + return -1; + } + UniqueFree SegsOwner(Segs); + + St = pQuerySegDescs(Segs, &NumSegs); + if (St != PROF_HSA_STATUS_SUCCESS) { + PROF_WARN("query_segment_descriptors(fetch) failed (0x%x)\n", St); + return -1; } - /* Warn about unprocessed TUs; skip cleared slots (already drained). */ - lockDynamicModules(); - for (int i = 0; i < NumDynamicModules; ++i) { - OffloadDynamicModuleInfo *MI = &DynamicModules[i]; - if (!MI->ModulePtr) + if (isVerboseMode()) + PROF_NOTE("query_segment_descriptors: %zu segments\n", NumSegs); + + /* Walk unique (agent, executable) pairs. */ + enum { kMaxPairs = 512 }; + uint64_t SeenAgents[kMaxPairs]; + uint64_t SeenExecs[kMaxPairs]; + int NumPairs = 0; + int IterFailures = 0; + + for (size_t i = 0; i < NumSegs; ++i) { + if (Segs[i].executable.handle == 0 || Segs[i].agent.handle == 0) continue; - for (int t = 0; t < MI->NumTUs; ++t) { - if (!MI->TUs[t].Processed) { - PROF_WARN("dynamic module %p TU %d was not processed before exit\n", - MI->ModulePtr, t); - Ret = -1; + + int Seen = 0; + for (int j = 0; j < NumPairs; ++j) + if (SeenAgents[j] == Segs[i].agent.handle && + SeenExecs[j] == Segs[i].executable.handle) { + Seen = 1; + break; } + if (Seen) + continue; + if (NumPairs < kMaxPairs) { + SeenAgents[NumPairs] = Segs[i].agent.handle; + SeenExecs[NumPairs] = Segs[i].executable.handle; + NumPairs++; } + + const char *Arch = nullptr; + for (int k = 0; k < W.num_agents; ++k) + if (W.agents[k].agent.handle == Segs[i].agent.handle) { + Arch = W.agents[k].arch; + break; + } + if (!Arch) + continue; /* not a GPU agent we collected */ + + SymbolState S; + __builtin_memset(&S, 0, sizeof(S)); + S.arch = Arch; + if (isVerboseMode()) + PROF_NOTE("walking executable 0x%llx on %s\n", + (unsigned long long)Segs[i].executable.handle, Arch); + prof_hsa_status_t IterSt = + pHsaExecIterAgentSyms(Segs[i].executable, Segs[i].agent, onSymbol, &S); + if (IterSt != PROF_HSA_STATUS_SUCCESS && + IterSt != PROF_HSA_STATUS_INFO_BREAK) { + PROF_WARN("hsa_executable_iterate_agent_symbols on executable 0x%llx " + "failed (0x%x)\n", + (unsigned long long)Segs[i].executable.handle, IterSt); + IterFailures++; + } + W.total_found += S.found; + W.total_drained += S.drained; } - unlockDynamicModules(); - if (Ret != 0) - PROF_WARN("%s\n", "failed to collect device profile data"); - return Ret; + if (isVerboseMode()) + PROF_NOTE("walk complete: agents=%d pairs=%d found=%d drained=%d " + "iter-failures=%d\n", + W.num_agents, NumPairs, W.total_found, W.total_drained, + IterFailures); + + if (W.total_found > 0 && W.total_drained == 0) { + PROF_WARN("found %d device profile symbol(s) but drained 0\n", + W.total_found); + return -1; + } + /* Latch only when we actually drained data. We deliberately do NOT latch the + * "walked everything but found no instrumented code object" case: that can + * happen on an early collect call (the host-write forwarder may run before any + * kernel launch, i.e. before atexit), and latching it would suppress the real + * atexit drain once kernels do run. Repeating a no-op walk is cheap and + * harmless, so leaving the latch clear here is the safe choice. */ + if (W.total_drained > 0) + __atomic_store_n(&DrainCompleted, 1, __ATOMIC_RELEASE); + return (IterFailures > 0) ? -1 : 0; } -/* Interceptors for hipModuleLoad* / hipModuleUnload. Linux only. */ - -#if defined(__linux__) && !defined(_WIN32) +/* -------------------------------------------------------------------------- */ +/* Public entry points */ +/* -------------------------------------------------------------------------- */ -INTERCEPTOR(int, hipModuleLoad, void **module, const char *fname) { - int rc = REAL(hipModuleLoad)(module, fname); - /* Pass NULL image: no in-memory ELF is available for filename loads, - * so the register hook skips symbol enumeration. */ - __llvm_profile_offload_register_dynamic_module(rc, module, nullptr); - return rc; +/* Called from the host write path (InstrProfilingFile.c) when the host TUs are + * instrumented. Independent of, and idempotent with, the atexit drain. */ +extern "C" int __llvm_profile_hip_collect_device_data(void) { + return drainDevices(); } -INTERCEPTOR(int, hipModuleLoadData, void **module, const void *image) { - int rc = REAL(hipModuleLoadData)(module, image); - __llvm_profile_offload_register_dynamic_module(rc, module, image); - return rc; -} +/* Legacy registration entry points from the previous host-shadow design, kept + * as no-ops so objects compiled against the old runtime still link. */ +extern "C" void __llvm_profile_offload_register_shadow_variable(void *) {} +extern "C" void +__llvm_profile_offload_register_section_shadow_variable(void *) {} +extern "C" void __llvm_profile_offload_register_dynamic_module(int, void **, + const void *) {} +extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *) {} -INTERCEPTOR(int, hipModuleLoadDataEx, void **module, const void *image, - unsigned numOptions, void **options, void **optionValues) { - int rc = REAL(hipModuleLoadDataEx)(module, image, numOptions, options, - optionValues); - __llvm_profile_offload_register_dynamic_module(rc, module, image); - return rc; -} +/* -------------------------------------------------------------------------- */ +/* Constructor */ +/* -------------------------------------------------------------------------- */ -INTERCEPTOR(int, hipModuleUnload, void *module) { - /* Drain counters before the module is destroyed; device addresses - * captured at register time are invalid after unload. */ - __llvm_profile_offload_unregister_dynamic_module(module); - return REAL(hipModuleUnload)(module); -} +static void atexitDrain(void) { (void)drainDevices(); } -__attribute__((constructor)) static void installHipModuleInterceptors() { - /* Skip when the HIP runtime is not loaded. INTERCEPT_FUNCTION uses the - * sanitizer interception framework, which can perturb dlsym/PLT state for - * the rest of the process even when the target symbol is absent; non-HIP - * programs linked with libclang_rt.profile.a must see zero side effects. */ - if (!dlsym(RTLD_DEFAULT, "hipModuleLoad")) - return; - if (!INTERCEPT_FUNCTION(hipModuleLoad)) - return; - if (isVerboseMode()) - PROF_NOTE("%s", "Installing hipModuleLoad*/hipModuleUnload interceptors\n"); - INTERCEPT_FUNCTION(hipModuleLoadData); - INTERCEPT_FUNCTION(hipModuleLoadDataEx); - INTERCEPT_FUNCTION(hipModuleUnload); +__attribute__((constructor)) static void profROCmInit(void) { + // Resolve and hsa_init now so HSA's atexit(hsa_shut_down) is registered + // before our atexit(drainDevices); LIFO then runs our drain while HSA is + // still alive. Failure here is non-fatal: a host-only program without ROCm + // simply gets no device drain. + (void)loadRuntimePointers(); + atexit(atexitDrain); } -#endif /* __linux__ */ +#endif // !defined(__NVPTX__) && !defined(__AMDGPU__) && !defined(_WIN32) diff --git a/compiler-rt/lib/profile/InstrProfilingPlatformROCmWindows.cpp b/compiler-rt/lib/profile/InstrProfilingPlatformROCmWindows.cpp new file mode 100644 index 0000000000000..b5ff949e2d3d4 --- /dev/null +++ b/compiler-rt/lib/profile/InstrProfilingPlatformROCmWindows.cpp @@ -0,0 +1,909 @@ +//===- InstrProfilingPlatformROCmWindows.cpp - ROCm platform (Windows) ---===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// Windows host-side ROCm/HIP device profile collection (legacy host-shadow + +// hipMemcpy drain). Windows has no HSA introspection, so it keeps this method; +// Linux uses HSA introspection in InstrProfilingPlatformROCm.cpp instead. The +// two files are mutually exclusive by platform guard so exactly one provides +// the __llvm_profile_offload_* / __llvm_profile_hip_collect_device_data symbols. +// +//===----------------------------------------------------------------------===// + +#if defined(_WIN32) + +extern "C" { +#include "InstrProfiling.h" +#include "InstrProfilingInternal.h" +#include "InstrProfilingPort.h" +} + +#include "interception/interception.h" +// C library headers (not etc.): clang_rt.profile is built with +// -nostdinc++ and avoids the C++ standard library (see profile/CMakeLists.txt). +#include +#include +#include +#include + +#ifdef _WIN32 +#define WIN32_LEAN_AND_MEAN +#include +#else +#include +#include +#endif + +/* Serialize one-time HIP loader resolution and DynamicModules mutations. + * Inline to avoid a sanitizer_common dependency. */ +#ifdef _WIN32 +static INIT_ONCE HipLoadedOnce = INIT_ONCE_STATIC_INIT; +static CRITICAL_SECTION DynamicModulesLock; +static INIT_ONCE DynamicModulesLockInit = INIT_ONCE_STATIC_INIT; +static BOOL CALLBACK initDynamicModulesLockCb(PINIT_ONCE, PVOID, PVOID *) { + InitializeCriticalSection(&DynamicModulesLock); + return TRUE; +} +static void lockDynamicModules(void) { + InitOnceExecuteOnce(&DynamicModulesLockInit, initDynamicModulesLockCb, NULL, + NULL); + EnterCriticalSection(&DynamicModulesLock); +} +static void unlockDynamicModules(void) { + LeaveCriticalSection(&DynamicModulesLock); +} +#else +static pthread_once_t HipLoadedOnce = PTHREAD_ONCE_INIT; +static pthread_mutex_t DynamicModulesLock = PTHREAD_MUTEX_INITIALIZER; +static void lockDynamicModules(void) { + pthread_mutex_lock(&DynamicModulesLock); +} +static void unlockDynamicModules(void) { + pthread_mutex_unlock(&DynamicModulesLock); +} +#endif + +static int processDeviceOffloadPrf(void *DeviceOffloadPrf, int TUIndex, + const char *Target); + +static int isVerboseMode() { + static int IsVerbose = -1; + if (IsVerbose == -1) + IsVerbose = getenv("LLVM_PROFILE_VERBOSE") != nullptr; + return IsVerbose; +} + +/* -------------------------------------------------------------------------- */ +/* Dynamic loading of HIP runtime symbols */ +/* -------------------------------------------------------------------------- */ + +typedef int (*hipGetSymbolAddressTy)(void **, const void *); +typedef int (*hipMemcpyTy)(void *, const void *, size_t, int); +typedef int (*hipModuleGetGlobalTy)(void **, size_t *, void *, const char *); +typedef int (*hipGetDeviceCountTy)(int *); +typedef int (*hipGetDeviceTy)(int *); +typedef int (*hipSetDeviceTy)(int); + +/* Minimal hipDeviceProp_t (HIP 6.x R0600): only gcnArchName at offset 1160 + * is read. Padded to 4096 to tolerate ABI growth. */ +typedef struct { + char padding[1160]; + char gcnArchName[256]; + char tail_padding[2680]; +} HipDevicePropMinimal; +typedef int (*hipGetDevicePropertiesTy)(HipDevicePropMinimal *, int); + +static hipGetSymbolAddressTy pHipGetSymbolAddress = nullptr; +static hipMemcpyTy pHipMemcpy = nullptr; +static hipModuleGetGlobalTy pHipModuleGetGlobal = nullptr; +static hipGetDeviceCountTy pHipGetDeviceCount = nullptr; +static hipGetDeviceTy pHipGetDevice = nullptr; +static hipSetDeviceTy pHipSetDevice = nullptr; +static hipGetDevicePropertiesTy pHipGetDeviceProperties = nullptr; + +static int NumDevices = 0; +/* 256 matches hipDeviceProp_t::gcnArchName, the source field width. */ +static char (*DeviceArchNames)[256] = nullptr; + +/* -------------------------------------------------------------------------- */ +/* Device-to-host copies */ +/* Keep HIP-only to avoid an HSA dependency. */ +/* -------------------------------------------------------------------------- */ + +static void doEnsureHipLoaded(void) { + if (!__interception::DynamicLoaderAvailable()) { + if (isVerboseMode()) + PROF_NOTE("%s", "Dynamic library loading not available - " + "HIP profiling disabled\n"); + return; + } + +#ifdef _WIN32 + static const char HipLibName[] = "amdhip64.dll"; +#else + static const char HipLibName[] = "libamdhip64.so"; +#endif + + void *Handle = __interception::OpenLibrary(HipLibName); + if (!Handle) + return; + + pHipGetSymbolAddress = (hipGetSymbolAddressTy)__interception::LookupSymbol( + Handle, "hipGetSymbolAddress"); + pHipMemcpy = (hipMemcpyTy)__interception::LookupSymbol(Handle, "hipMemcpy"); + pHipModuleGetGlobal = (hipModuleGetGlobalTy)__interception::LookupSymbol( + Handle, "hipModuleGetGlobal"); + pHipGetDeviceCount = (hipGetDeviceCountTy)__interception::LookupSymbol( + Handle, "hipGetDeviceCount"); + pHipGetDevice = + (hipGetDeviceTy)__interception::LookupSymbol(Handle, "hipGetDevice"); + pHipSetDevice = + (hipSetDeviceTy)__interception::LookupSymbol(Handle, "hipSetDevice"); + pHipGetDeviceProperties = + (hipGetDevicePropertiesTy)__interception::LookupSymbol( + Handle, "hipGetDevicePropertiesR0600"); + if (!pHipGetDeviceProperties) + pHipGetDeviceProperties = + (hipGetDevicePropertiesTy)__interception::LookupSymbol( + Handle, "hipGetDeviceProperties"); + + if (pHipGetDeviceCount && pHipGetDeviceProperties) { + int Count = 0; + if (pHipGetDeviceCount(&Count) == 0 && Count > 0) { + DeviceArchNames = (char (*)[256])calloc(Count, sizeof(*DeviceArchNames)); + if (!DeviceArchNames) { + PROF_ERR("%s\n", "failed to allocate device arch name table"); + return; + } + HipDevicePropMinimal Prop; + for (int i = 0; i < Count; ++i) { + __builtin_memset(&Prop, 0, sizeof(Prop)); + if (pHipGetDeviceProperties(&Prop, i) == 0) { + strncpy(DeviceArchNames[i], Prop.gcnArchName, + sizeof(DeviceArchNames[i]) - 1); + DeviceArchNames[i][sizeof(DeviceArchNames[i]) - 1] = '\0'; + if (isVerboseMode()) + PROF_NOTE("Device %d arch: %s\n", i, DeviceArchNames[i]); + } + } + NumDevices = Count; + } + } +} + +#ifdef _WIN32 +static BOOL CALLBACK ensureHipLoadedCb(PINIT_ONCE, PVOID, PVOID *) { + doEnsureHipLoaded(); + return TRUE; +} +#endif + +static void ensureHipLoaded(void) { +#ifdef _WIN32 + InitOnceExecuteOnce(&HipLoadedOnce, ensureHipLoadedCb, NULL, NULL); +#else + pthread_once(&HipLoadedOnce, doEnsureHipLoaded); +#endif +} + +/* -------------------------------------------------------------------------- */ +/* Public wrappers that forward to the loaded HIP symbols */ +/* -------------------------------------------------------------------------- */ + +static int hipGetSymbolAddress(void **devPtr, const void *symbol) { + ensureHipLoaded(); + return pHipGetSymbolAddress ? pHipGetSymbolAddress(devPtr, symbol) : -1; +} + +static int hipMemcpy(void *dest, const void *src, size_t len, + int kind /*2=DToH*/) { + ensureHipLoaded(); + return pHipMemcpy ? pHipMemcpy(dest, src, len, kind) : -1; +} + +/* Device section symbols must be registered with CLR first; otherwise + * hipMemcpy may take a CPU path and crash. */ +static int memcpyDeviceToHost(void *Dst, const void *Src, size_t Size) { + return hipMemcpy(Dst, Src, Size, 2 /* DToH */); +} + +static int hipModuleGetGlobal(void **DevPtr, size_t *Bytes, void *Module, + const char *Name) { + ensureHipLoaded(); + return pHipModuleGetGlobal ? pHipModuleGetGlobal(DevPtr, Bytes, Module, Name) + : -1; +} + +static int hipGetDevice(int *DeviceId) { + ensureHipLoaded(); + return pHipGetDevice ? pHipGetDevice(DeviceId) : -1; +} + +static int hipSetDevice(int DeviceId) { + ensureHipLoaded(); + return pHipSetDevice ? pHipSetDevice(DeviceId) : -1; +} + +static const char *getDeviceArchName(int DeviceId) { + if (DeviceId < 0 || DeviceId >= NumDevices || !DeviceArchNames[DeviceId][0]) + return "amdgpu"; + return DeviceArchNames[DeviceId]; +} + +/* -------------------------------------------------------------------------- */ +/* Dynamic module tracking */ +/* -------------------------------------------------------------------------- */ + +/* Per-TU profile entry inside a dynamic module. + * A single dynamic module may contain multiple TUs (e.g. -fgpu-rdc). */ +typedef struct { + void *DeviceVar; /* device address of __llvm_profile_sections_ */ + int Processed; /* 0 = not yet collected, 1 = data already copied */ +} OffloadDynamicTUInfo; + +/* One entry per hipModuleLoad call. */ +typedef struct { + void *ModulePtr; /* hipModule_t handle */ + OffloadDynamicTUInfo *TUs; /* array of per-TU entries */ + int NumTUs; + int CapTUs; +} OffloadDynamicModuleInfo; + +static OffloadDynamicModuleInfo *DynamicModules = nullptr; +static int NumDynamicModules = 0; +static int CapDynamicModules = 0; + +/* -------------------------------------------------------------------------- */ +/* ELF symbol enumeration (manual parse: compiler-rt cannot link LLVM Support) + */ +/* -------------------------------------------------------------------------- */ + +#if __has_include() +#include + +/* Callback invoked for every matching symbol name found in the ELF image. + * Return 0 to continue iteration, non-zero to stop. */ +typedef int (*SymbolCallback)(const char *Name, void *UserData); + +/* If Image is a clang offload bundle, return a pointer to the first embedded + * ELF. Returns Image if not a bundle, nullptr if a bundle holds no ELF. */ +static const void *unwrapOffloadBundle(const void *Image) { + static const char BundleMagic[] = "__CLANG_OFFLOAD_BUNDLE__"; + if (memcmp(Image, BundleMagic, sizeof(BundleMagic) - 1) != 0) + return Image; /* Not a bundle, return as-is. */ + + const char *Buf = (const char *)Image; + uint64_t NumEntries; + __builtin_memcpy(&NumEntries, Buf + sizeof(BundleMagic) - 1, + sizeof(uint64_t)); + + /* Walk the entry table (starts at offset 32). */ + const char *Cursor = Buf + 32; + for (uint64_t I = 0; I < NumEntries; ++I) { + uint64_t EntryOffset, EntrySize, IDSize; + __builtin_memcpy(&EntryOffset, Cursor, sizeof(EntryOffset)); + Cursor += sizeof(EntryOffset); + __builtin_memcpy(&EntrySize, Cursor, sizeof(EntrySize)); + Cursor += sizeof(EntrySize); + __builtin_memcpy(&IDSize, Cursor, sizeof(IDSize)); + Cursor += sizeof(IDSize); + Cursor += IDSize; /* skip entry ID */ + + if (EntrySize >= sizeof(Elf64_Ehdr)) { + const Elf64_Ehdr *E = (const Elf64_Ehdr *)(Buf + EntryOffset); + if (E->e_ident[EI_MAG0] == ELFMAG0 && E->e_ident[EI_MAG1] == ELFMAG1 && + E->e_ident[EI_MAG2] == ELFMAG2 && E->e_ident[EI_MAG3] == ELFMAG3) { + return (const void *)(Buf + EntryOffset); + } + } + } + + PROF_WARN("%s", "offload bundle contains no valid ELF entries\n"); + return nullptr; +} + +/* Invoke CB for every global symbol in Image (an AMDGPU ELF or offload bundle) + * whose name starts with PREFIX. Image may be null. */ +static void enumerateElfSymbols(const void *Image, const char *Prefix, + SymbolCallback CB, void *UserData) { + if (!Image) + return; + + Image = unwrapOffloadBundle(Image); + if (!Image) + return; + + const Elf64_Ehdr *Ehdr = (const Elf64_Ehdr *)Image; + if (Ehdr->e_ident[EI_MAG0] != ELFMAG0 || Ehdr->e_ident[EI_MAG1] != ELFMAG1 || + Ehdr->e_ident[EI_MAG2] != ELFMAG2 || Ehdr->e_ident[EI_MAG3] != ELFMAG3) { + if (isVerboseMode()) + PROF_NOTE("%s", "Image is not a valid ELF, skipping enumeration\n"); + return; + } + + size_t PrefixLen = strlen(Prefix); + const char *Base = (const char *)Image; + const Elf64_Shdr *Shdrs = (const Elf64_Shdr *)(Base + Ehdr->e_shoff); + + for (int i = 0; i < Ehdr->e_shnum; ++i) { + if (Shdrs[i].sh_type != SHT_SYMTAB) + continue; + + const Elf64_Sym *Syms = (const Elf64_Sym *)(Base + Shdrs[i].sh_offset); + int NumSyms = Shdrs[i].sh_size / sizeof(Elf64_Sym); + /* String table is the section referenced by sh_link. */ + const char *StrTab = Base + Shdrs[Shdrs[i].sh_link].sh_offset; + + for (int j = 0; j < NumSyms; ++j) { + if (Syms[j].st_name == 0) + continue; + const char *Name = StrTab + Syms[j].st_name; + if (strncmp(Name, Prefix, PrefixLen) == 0) { + if (CB(Name, UserData)) + return; + } + } + } +} + +/* State passed through the enumeration callback. */ +typedef struct { + void *Module; /* hipModule_t */ + OffloadDynamicModuleInfo *ModInfo; +} EnumState; + +/* Register one __llvm_profile_sections_ symbol on the module entry. + * hipModuleGetGlobal also registers the device address with CLR so hipMemcpy + * can copy from it later. */ +static int registerPrfSymbol(const char *Name, void *UserData) { + EnumState *S = (EnumState *)UserData; + OffloadDynamicModuleInfo *MI = S->ModInfo; + + /* The symbol is the per-TU sections struct itself, not a pointer + * indirection, so this address is the hipMemcpy source. */ + void *DeviceVar = nullptr; + size_t Bytes = 0; + if (hipModuleGetGlobal(&DeviceVar, &Bytes, S->Module, Name) != 0) { + PROF_WARN("failed to get symbol %s for module %p\n", Name, S->Module); + return 0; /* continue */ + } + + if (MI->NumTUs >= MI->CapTUs) { + int NewCap = MI->CapTUs ? MI->CapTUs * 2 : 4; + OffloadDynamicTUInfo *New = (OffloadDynamicTUInfo *)realloc( + MI->TUs, NewCap * sizeof(OffloadDynamicTUInfo)); + if (!New) { + PROF_ERR("%s\n", "failed to grow TU array"); + return 0; + } + MI->TUs = New; + MI->CapTUs = NewCap; + } + OffloadDynamicTUInfo *TU = &MI->TUs[MI->NumTUs++]; + TU->DeviceVar = DeviceVar; + TU->Processed = 0; + + (void)Name; + return 0; /* continue enumeration */ +} + +#endif /* __has_include() */ + +/* -------------------------------------------------------------------------- */ +/* Registration / un-registration helpers */ +/* -------------------------------------------------------------------------- */ + +extern "C" void +__llvm_profile_offload_register_dynamic_module(int ModuleLoadRc, void **Ptr, + const void *Image) { + if (ModuleLoadRc) + return; + + lockDynamicModules(); + + if (isVerboseMode()) + PROF_NOTE("Registering loaded module %d: rc=%d, module=%p, image=%p\n", + NumDynamicModules, ModuleLoadRc, *Ptr, Image); + + if (NumDynamicModules >= CapDynamicModules) { + int NewCap = CapDynamicModules ? CapDynamicModules * 2 : 64; + OffloadDynamicModuleInfo *New = (OffloadDynamicModuleInfo *)realloc( + DynamicModules, NewCap * sizeof(OffloadDynamicModuleInfo)); + if (!New) { + unlockDynamicModules(); + return; + } + DynamicModules = New; + CapDynamicModules = NewCap; + } + + OffloadDynamicModuleInfo *MI = &DynamicModules[NumDynamicModules++]; + MI->ModulePtr = *Ptr; + MI->TUs = nullptr; + MI->NumTUs = 0; + MI->CapTUs = 0; + + /* Dynamic-module profiling needs ELF parsing for symbol enumeration. */ +#if __has_include() + EnumState State = {*Ptr, MI}; + enumerateElfSymbols(Image, "__llvm_profile_sections_", 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_profile_sections_* symbols found in module %p\n", + *Ptr); + } else if (isVerboseMode()) { + PROF_NOTE("Module %p: registered %d TU(s)\n", *Ptr, MI->NumTUs); + } + + unlockDynamicModules(); +} + +extern "C" void __llvm_profile_offload_unregister_dynamic_module(void *Ptr) { + lockDynamicModules(); + for (int i = 0; i < NumDynamicModules; ++i) { + OffloadDynamicModuleInfo *MI = &DynamicModules[i]; + + /* HIP recycles hipModule_t addresses; drained slots are cleared so a + * recycled handle finds the new slot, not the dead one. */ + if (MI->ModulePtr != Ptr) + continue; + + if (isVerboseMode()) + PROF_NOTE("Unregistering module %p (%d TUs)\n", MI->ModulePtr, + MI->NumTUs); + + static int NextTUIndex = 0; + 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; + } + int TUIndex = __atomic_fetch_add(&NextTUIndex, 1, __ATOMIC_RELAXED); + if (TU->DeviceVar) { + int CurDev = 0; + hipGetDevice(&CurDev); + const char *ArchName = getDeviceArchName(CurDev); + /* Encode TUIndex in Target so each drain writes a distinct profraw; + * otherwise back-to-back drains overwrite the same file. */ + char TargetWithTU[64]; + snprintf(TargetWithTU, sizeof(TargetWithTU), "%s.%d", ArchName, + TUIndex); + if (processDeviceOffloadPrf(TU->DeviceVar, TUIndex, TargetWithTU) == 0) + TU->Processed = 1; + else + PROF_WARN("failed to process profile data for module %p TU %d\n", Ptr, + t); + } + } + MI->ModulePtr = nullptr; + unlockDynamicModules(); + return; + } + + if (isVerboseMode()) + PROF_WARN("unregister called for unknown module %p\n", Ptr); + unlockDynamicModules(); +} + +/* Grow a void* array, doubling capacity (or starting at InitCap). */ +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; +} + +namespace { + +// free()-based scope guard. Use .release() to transfer ownership. +struct UniqueFree { + void *Ptr; + explicit UniqueFree(void *P = nullptr) : Ptr(P) {} + ~UniqueFree() { free(Ptr); } + UniqueFree(const UniqueFree &) = delete; + UniqueFree &operator=(const UniqueFree &) = delete; + char *get() const { return static_cast(Ptr); } + void reset(void *P) { + free(Ptr); + Ptr = P; + } + void *release() { + void *P = Ptr; + Ptr = nullptr; + return P; + } +}; + +} // 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; + + // Owns freshly malloc'd buffers; release() transfers ownership to the cache. + UniqueFree CntsOwner, DataOwner, NamesOwner; + + 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); + CntsOwner.reset(HostCountersBegin); + } + + 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); + DataOwner.reset(HostDataBegin); + } + + 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); + NamesOwner.reset(HostNamesBegin); + } + + if ((DataSize > 0 && !HostDataBegin) || + (CountersSize > 0 && !HostCountersBegin) || + (NamesSize > 0 && !HostNamesBegin)) { + PROF_ERR("%s\n", "failed to allocate host memory for device sections"); + return -1; + } + + 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; + } + + /* Cache buffers so RDC-mode multi-shadow drains can reuse them. + * release() prevents the scope guards from freeing what the cache owns. */ + if (!CntsReused && CountersSize > 0) { + CachedDevCntsBegin = DevCntsBegin; + CachedHostCnts = HostCountersBegin; + CachedCntsSize = CountersSize; + CntsOwner.release(); + } + if (!DataReused && DataSize > 0) { + CachedDevDataBegin = DevDataBegin; + CachedHostData = HostDataBegin; + CachedDataSize = DataSize; + DataOwner.release(); + } + if (!NamesReused && NamesSize > 0) { + CachedDevNamesBegin = DevNamesBegin; + CachedHostNames = HostNamesBegin; + CachedNamesSize = NamesSize; + NamesOwner.release(); + } + + 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; + UniqueFree ContiguousBuf(malloc(ContiguousBufferSize)); + if (!ContiguousBuf.get()) { + PROF_ERR("%s\n", "failed to allocate contiguous buffer"); + return -1; + } + char *ContiguousBuffer = ContiguousBuf.get(); + __builtin_memset(ContiguousBuffer, 0, ContiguousBufferSize); + + char *BufDataBegin = ContiguousBuffer; + char *BufCountersBegin = + ContiguousBuffer + DataSize + PaddingBytesBeforeCounters; + char *BufNamesBegin = BufCountersBegin + CountersSize; + + __builtin_memcpy(BufDataBegin, HostDataBegin, DataSize); + __builtin_memcpy(BufCountersBegin, HostCountersBegin, CountersSize); + __builtin_memcpy(BufNamesBegin, HostNamesBegin, NamesSize); + + // CounterPtr is a device-relative offset; relocate it for the file layout + // where the Data section precedes Counters. + __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)); + __builtin_memcpy((char *)RelocatedData + i * sizeof(__llvm_profile_data) + + offsetof(__llvm_profile_data, CounterPtr), + &NewRelativeOffset, sizeof(NewRelativeOffset)); + } + __builtin_memset((char *)RelocatedData + i * sizeof(__llvm_profile_data) + + offsetof(__llvm_profile_data, BitmapPtr), + 0, + sizeof(RelocatedData[i].BitmapPtr) + + sizeof(RelocatedData[i].FunctionPointer) + + sizeof(RelocatedData[i].Values)); + } + + /* Target already encodes TUIndex when needed. */ + (void)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 *DeviceSections = nullptr; + if (hipGetSymbolAddress(&DeviceSections, ShadowVar) != 0) { + PROF_WARN("failed to get symbol address for shadow variable %p\n", + ShadowVar); + return -1; + } + /* DeviceSections points at the per-TU sections struct itself. */ + return processDeviceOffloadPrf(DeviceSections, TUIndex, Target); +} + +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): drain from every device. */ + 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) { + /* RDC-mode multi-shadow drains need a distinct profraw per TU; + * single-TU programs keep the bare arch target. */ + const char *Target = ArchName; + char TargetWithIdx[64]; + if (NumShadowVariables > 1) { + snprintf(TargetWithIdx, sizeof(TargetWithIdx), "%s.%d", ArchName, i); + Target = TargetWithIdx; + } + if (processShadowVariable(OffloadShadowVariables[i], i, Target) != 0) + Ret = -1; + } + } + + if (OrigDevice >= 0) + hipSetDevice(OrigDevice); + } + + /* Warn about unprocessed TUs; skip cleared slots (already drained). */ + lockDynamicModules(); + for (int i = 0; i < NumDynamicModules; ++i) { + OffloadDynamicModuleInfo *MI = &DynamicModules[i]; + if (!MI->ModulePtr) + continue; + for (int t = 0; t < MI->NumTUs; ++t) { + if (!MI->TUs[t].Processed) { + PROF_WARN("dynamic module %p TU %d was not processed before exit\n", + MI->ModulePtr, t); + Ret = -1; + } + } + } + unlockDynamicModules(); + + if (Ret != 0) + PROF_WARN("%s\n", "failed to collect device profile data"); + return Ret; +} + +/* Interceptors for hipModuleLoad* / hipModuleUnload. Linux only. */ + +#if defined(__linux__) && !defined(_WIN32) + +INTERCEPTOR(int, hipModuleLoad, void **module, const char *fname) { + int rc = REAL(hipModuleLoad)(module, fname); + /* Pass NULL image: no in-memory ELF is available for filename loads, + * so the register hook skips symbol enumeration. */ + __llvm_profile_offload_register_dynamic_module(rc, module, nullptr); + return rc; +} + +INTERCEPTOR(int, hipModuleLoadData, void **module, const void *image) { + int rc = REAL(hipModuleLoadData)(module, image); + __llvm_profile_offload_register_dynamic_module(rc, module, image); + return rc; +} + +INTERCEPTOR(int, hipModuleLoadDataEx, void **module, const void *image, + unsigned numOptions, void **options, void **optionValues) { + int rc = REAL(hipModuleLoadDataEx)(module, image, numOptions, options, + optionValues); + __llvm_profile_offload_register_dynamic_module(rc, module, image); + return rc; +} + +INTERCEPTOR(int, hipModuleUnload, void *module) { + /* Drain counters before the module is destroyed; device addresses + * captured at register time are invalid after unload. */ + __llvm_profile_offload_unregister_dynamic_module(module); + return REAL(hipModuleUnload)(module); +} + +__attribute__((constructor)) static void installHipModuleInterceptors() { + /* Skip when the HIP runtime is not loaded. INTERCEPT_FUNCTION uses the + * sanitizer interception framework, which can perturb dlsym/PLT state for + * the rest of the process even when the target symbol is absent; non-HIP + * programs linked with libclang_rt.profile.a must see zero side effects. */ + if (!dlsym(RTLD_DEFAULT, "hipModuleLoad")) + return; + if (!INTERCEPT_FUNCTION(hipModuleLoad)) + return; + if (isVerboseMode()) + PROF_NOTE("%s", "Installing hipModuleLoad*/hipModuleUnload interceptors\n"); + INTERCEPT_FUNCTION(hipModuleLoadData); + INTERCEPT_FUNCTION(hipModuleLoadDataEx); + INTERCEPT_FUNCTION(hipModuleUnload); +} + +#endif /* __linux__ */ + +#endif /* defined(_WIN32) */ diff --git a/compiler-rt/test/profile/AMDGPU/device-basic.hip b/compiler-rt/test/profile/AMDGPU/device-basic.hip new file mode 100644 index 0000000000000..4fcf044802240 --- /dev/null +++ b/compiler-rt/test/profile/AMDGPU/device-basic.hip @@ -0,0 +1,67 @@ +// Basic HIP device PGO drain end-to-end: a host + device .profraw are written +// at exit (the device one arch-prefixed), they merge, the merged profile +// contains the device kernel's counters, and llvm-cov reports device-side +// coverage. Covers both non-RDC and RDC device compiles. +// +// REQUIRES: hip, amdgpu + +// RUN: rm -rf %t.dir && mkdir -p %t.dir + +// --- non-RDC --- +// RUN: %clang -x hip --offload-arch=%amdgpu_arch -fno-gpu-rdc \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/a.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// RUN: env LLVM_PROFILE_FILE=%t.dir/host.%%p.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: %t.dir/a.out +// A device profraw (arch-prefixed) must have been drained alongside the host one. +// RUN: ls %t.dir/gfx*.profraw +// RUN: llvm-profdata merge %t.dir/*.profraw -o %t.dir/a.profdata +// RUN: llvm-profdata show --all-functions %t.dir/a.profdata \ +// RUN: | FileCheck --check-prefix=FUNCS %s +// Confirm the embedded device image is extractable (failure here is the real +// cause of any downstream llvm-cov failure, so let it propagate). +// RUN: llvm-objdump --offloading %t.dir/a.out > /dev/null +// RUN: llvm-cov report %t.dir/a.out.0.hip-amdgcn-amd-amdhsa--*gfx* \ +// RUN: -instr-profile=%t.dir/a.profdata 2>&1 | FileCheck --check-prefix=COV %s + +// --- RDC --- +// RUN: rm -f %t.dir/*.profraw +// RUN: %clang -x hip --offload-arch=%amdgpu_arch -fgpu-rdc \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/b.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// RUN: env LLVM_PROFILE_FILE=%t.dir/host.%%p.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: %t.dir/b.out +// RUN: ls %t.dir/gfx*.profraw +// RUN: llvm-profdata merge %t.dir/*.profraw -o %t.dir/b.profdata +// RUN: llvm-profdata show --all-functions %t.dir/b.profdata \ +// RUN: | FileCheck --check-prefix=FUNCS %s + +#include + +__global__ void addk(int *p) { + if (*p > 0) + *p += 1; + else + *p -= 1; +} + +int main() { + int *d = nullptr; + if (hipMalloc(&d, sizeof(int)) != hipSuccess) + return 2; + int h = 5; + (void)hipMemcpy(d, &h, sizeof(int), hipMemcpyHostToDevice); + addk<<<1, 1>>>(d); + (void)hipMemcpy(&h, d, sizeof(int), hipMemcpyDeviceToHost); + (void)hipFree(d); + return h > 0 ? 0 : 1; +} + +// The merged profile contains both the host main and the device kernel, +// proving the device counters were drained and merged. +// FUNCS-DAG: addk +// FUNCS-DAG: main + +// COV: TOTAL diff --git a/compiler-rt/test/profile/AMDGPU/device-early-collect.hip b/compiler-rt/test/profile/AMDGPU/device-early-collect.hip new file mode 100644 index 0000000000000..3e2c6e84e26c2 --- /dev/null +++ b/compiler-rt/test/profile/AMDGPU/device-early-collect.hip @@ -0,0 +1,68 @@ +// M1 regression: calling __llvm_profile_hip_collect_device_data() before any +// kernel has been launched must not poison the later atexit drain. The early +// call sees "no instrumented code object loaded yet" (a transient no-op) and +// must not latch the drain as completed; otherwise the post-launch atexit +// pass produces no device .profraw and we silently lose device counters. +// +// REQUIRES: hip, amdgpu +// Guards the Linux introspection drain's DrainCompleted latch; the Windows +// host-shadow drain has no such latch (it tracks per-TU Processed flags). +// UNSUPPORTED: windows + +// RUN: rm -rf %t.dir && mkdir -p %t.dir +// RUN: %clang -x hip --offload-arch=%amdgpu_arch -fno-gpu-rdc \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/a.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// RUN: env LLVM_PROFILE_FILE=%t.dir/host.%%p.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: %t.dir/a.out +// Both the host profraw and at least one device profraw (gfx-prefixed) must +// have been produced, despite the early collection attempt. +// RUN: ls %t.dir/host.*.profraw +// RUN: ls %t.dir/gfx*.profraw +// And the merged profile must contain the device kernel that was launched +// *after* the early collect. +// RUN: llvm-profdata merge %t.dir/*.profraw -o %t.dir/a.profdata +// RUN: llvm-profdata show --all-functions %t.dir/a.profdata \ +// RUN: | FileCheck %s + +#include + +// Declared by libclang_rt.profile-.a; we call it directly to +// simulate any caller that drains device counters at an arbitrary point in +// the program lifetime (e.g. a per-iteration profile dump). +extern "C" int __llvm_profile_hip_collect_device_data(void); + +__global__ void post_collect_kernel(int *p) { + if (*p > 0) + *p += 1; + else + *p -= 1; +} + +int main() { + // (1) Early collection -- runs before any kernel launch. The drainer + // finds either no GPU agents, no loaded segments, or no instrumented + // bounds table, and returns 0 without latching DrainCompleted. + (void)__llvm_profile_hip_collect_device_data(); + + // (2) Now launch a kernel. HIP loads the device code object that carries + // the __llvm_profile_sections bounds table, executes our kernel, and + // populates the device-side counters. + int *d = nullptr; + if (hipMalloc(&d, sizeof(int)) != hipSuccess) + return 2; + int h = 5; + (void)hipMemcpy(d, &h, sizeof(int), hipMemcpyHostToDevice); + post_collect_kernel<<<1, 1>>>(d); + (void)hipMemcpy(&h, d, sizeof(int), hipMemcpyDeviceToHost); + (void)hipFree(d); + + // (3) Exit normally. The atexit drain runs and -- because step (1) did + // not latch DrainCompleted -- it walks the (now loaded) code object, + // finds __llvm_profile_sections, and emits the device .profraw. + return h > 0 ? 0 : 1; +} + +// CHECK-DAG: post_collect_kernel +// CHECK-DAG: main diff --git a/compiler-rt/test/profile/AMDGPU/device-no-kernel.hip b/compiler-rt/test/profile/AMDGPU/device-no-kernel.hip new file mode 100644 index 0000000000000..a154308d725d8 --- /dev/null +++ b/compiler-rt/test/profile/AMDGPU/device-no-kernel.hip @@ -0,0 +1,44 @@ +// Independence / robustness: an instrumented HIP program that never launches a +// kernel still writes its host .profraw, and the device drain is a clean no-op +// (no crash, no spurious device .profraw). We assert the no-op condition +// directly via the runtime's verbose log rather than rely on HIP lazy-loading +// to leave the device code object unloaded -- the loader may load it for +// other reasons (e.g. eager registration), and in that case the drain +// legitimately walks it and reports zero instrumented sections / zero +// drained. Either outcome is correct. +// +// REQUIRES: hip, amdgpu +// The terminal conditions checked below ("no GPU agents", "no loaded +// segments", "drained=0") are Linux HSA-drain strings with no Windows analog. +// UNSUPPORTED: windows + +// RUN: rm -rf %t.dir && mkdir -p %t.dir +// RUN: %clang -x hip --offload-arch=%amdgpu_arch \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/a.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// RUN: env LLVM_PROFILE_FILE=%t.dir/host.%%p.profraw \ +// RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ +// RUN: LLVM_PROFILE_VERBOSE=1 \ +// RUN: %t.dir/a.out 2> %t.dir/verbose.log +// RUN: ls %t.dir/host.*.profraw +// No arch-prefixed device profraw should have been produced. +// RUN: not ls %t.dir/gfx*.profraw +// The drain must have run; one of these three terminal conditions must hold: +// - no GPU agents enumerated (test host has /dev/kfd but no usable agent) +// - no loaded code object segments at exit +// - the walk completed and drained=0 (no instrumented kernel was launched +// so the device code object either wasn't loaded or its bounds were +// empty/already drained) +// RUN: FileCheck --input-file=%t.dir/verbose.log %s +// CHECK: {{no GPU agents present|no loaded segments|drained=0}} + +#include + +// Defined but never launched. +__global__ void unused(int *p) { *p += 1; } + +int main() { + int n = 0; + (void)hipGetDeviceCount(&n); + return 0; +} diff --git a/compiler-rt/test/profile/AMDGPU/device-symbols.hip b/compiler-rt/test/profile/AMDGPU/device-symbols.hip new file mode 100644 index 0000000000000..f12283b7da636 --- /dev/null +++ b/compiler-rt/test/profile/AMDGPU/device-symbols.hip @@ -0,0 +1,42 @@ +// The decoupled drain reads only the canonical __llvm_profile_sections bounds +// table provided by the device profile runtime (InstrProfilingPlatformGPU.c), +// since clang no longer emits a per-TU struct. Assert that symbol is present +// in the device ELF's dynamic symbol table (protected visibility) for both +// non-RDC and RDC device compiles. This is the contract the drainer depends on. +// +// REQUIRES: hip, amdgpu + +// RUN: rm -rf %t.dir && mkdir -p %t.dir + +// --- non-RDC --- +// RUN: %clang -x hip --offload-arch=%amdgpu_arch -fno-gpu-rdc \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/a.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// Extraction failure here would make the readelf invocation succeed against +// an empty/missing file; surface it instead of hiding it behind `|| true`. +// RUN: llvm-objdump --offloading %t.dir/a.out > /dev/null +// RUN: llvm-readelf --dyn-syms %t.dir/a.out.0.hip-amdgcn-amd-amdhsa--*gfx* \ +// RUN: | FileCheck %s + +// --- RDC --- +// RUN: %clang -x hip --offload-arch=%amdgpu_arch -fgpu-rdc \ +// RUN: -fprofile-instr-generate -fcoverage-mapping %s -o %t.dir/b.out \ +// RUN: -L%hip_lib_path -lamdhip64 +// RUN: llvm-objdump --offloading %t.dir/b.out > /dev/null +// RUN: llvm-readelf --dyn-syms %t.dir/b.out.0.hip-amdgcn-amd-amdhsa--*gfx* \ +// RUN: | FileCheck %s + +// CHECK: PROTECTED {{.*}} __llvm_profile_sections + +#include + +__global__ void k(int *p) { *p += 1; } + +int main() { + int *d = nullptr; + if (hipMalloc(&d, sizeof(int)) != hipSuccess) + return 2; + k<<<1, 1>>>(d); + (void)hipFree(d); + return 0; +} diff --git a/compiler-rt/test/profile/AMDGPU/lit.local.cfg.py b/compiler-rt/test/profile/AMDGPU/lit.local.cfg.py new file mode 100644 index 0000000000000..5148dd6b9e2f2 --- /dev/null +++ b/compiler-rt/test/profile/AMDGPU/lit.local.cfg.py @@ -0,0 +1,4 @@ +# Device-profile drain tests: require an AMD GPU (and, implicitly, the amdgcn +# device profile runtime in the resource directory and a ROCm/HIP install). +if "amdgpu" not in config.available_features: + config.unsupported = True diff --git a/compiler-rt/test/profile/GPU/instrprof-hip-coverage.hip b/compiler-rt/test/profile/GPU/instrprof-hip-coverage.hip index 02dffc0c0fa15..a867c30f0edfb 100644 --- a/compiler-rt/test/profile/GPU/instrprof-hip-coverage.hip +++ b/compiler-rt/test/profile/GPU/instrprof-hip-coverage.hip @@ -12,7 +12,9 @@ // RUN: | FileCheck %s --check-prefix=REPORT // // REPORT: instrprof-hip-coverage.hip -// REPORT-NOT: 0.00% +// No coverage column should be fully uncovered. Anchor on a non-digit before +// the "0.00%" so this does not spuriously match e.g. "80.00%". +// REPORT-NOT: {{[^.0-9]0[.]00%}} #include #include diff --git a/compiler-rt/test/profile/GPU/instrprof-hip-multi-gpu.hip b/compiler-rt/test/profile/GPU/instrprof-hip-multi-gpu.hip index 740aaa326e60d..054ce3777bd1d 100644 --- a/compiler-rt/test/profile/GPU/instrprof-hip-multi-gpu.hip +++ b/compiler-rt/test/profile/GPU/instrprof-hip-multi-gpu.hip @@ -1,17 +1,23 @@ -// Test that HIP PGO doesn't crash on multi-GPU systems. -// The profile runtime should only collect from the active device at exit, -// not iterate all devices (which can trigger comgr lazy-load crashes). +// Test that HIP PGO works on multi-GPU systems. The HSA-introspection drain +// walks every GPU agent at exit and collects device counters from each loaded +// code object; it must do so without crashing (the previous host-shadow design +// relied on hipGetSymbolAddress and could hit a comgr lazy-load null deref when +// more than one device was present). // // REQUIRES: hip, amdgpu +// The HSA agent-walk drain is Linux-only; the Windows host-shadow drain +// collects only the current device and emits no "walk complete" line. +// UNSUPPORTED: windows // RUN: %clang -x hip -fprofile-instr-generate -fcoverage-mapping \ // RUN: --offload-arch=%amdgpu_arch %s -o %t -L%hip_lib_path -lamdhip64 // RUN: env LLVM_PROFILE_FILE=%t.profraw \ // RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ // RUN: LLVM_PROFILE_VERBOSE=1 %run %t 2>&1 | FileCheck %s // -// Verify device 0 data was collected and program didn't crash. -// CHECK: Collecting static profile data from device 0 -// CHECK: Successfully wrote device profile using shared API +// Verify the drain walked the GPU agents, copied device counters, and the +// program didn't crash on a multi-GPU system. +// CHECK: Copied device sections: +// CHECK: walk complete: agents={{[0-9]+}} pairs={{[0-9]+}} found={{[0-9]+}} drained={{[1-9][0-9]*}} // CHECK: PASS #include diff --git a/compiler-rt/test/profile/GPU/instrprof-hip-multiple-kernels.hip b/compiler-rt/test/profile/GPU/instrprof-hip-multiple-kernels.hip index ce1f5d7346ede..0fd6185b82441 100644 --- a/compiler-rt/test/profile/GPU/instrprof-hip-multiple-kernels.hip +++ b/compiler-rt/test/profile/GPU/instrprof-hip-multiple-kernels.hip @@ -14,8 +14,8 @@ // // All three kernels plus main should be profiled. // PROF-DAG: _Z4fillPii -// PROF-DAG: _Z5scalePixi -// PROF-DAG: _Z6negatePin +// PROF-DAG: _Z5scalePii +// PROF-DAG: _Z6negatePii // PROF-DAG: main // PROF: Total functions: 4 diff --git a/compiler-rt/test/profile/GPU/instrprof-hip-nondefault-device.hip b/compiler-rt/test/profile/GPU/instrprof-hip-nondefault-device.hip index 71501925e72e3..37bff9ecaa879 100644 --- a/compiler-rt/test/profile/GPU/instrprof-hip-nondefault-device.hip +++ b/compiler-rt/test/profile/GPU/instrprof-hip-nondefault-device.hip @@ -1,17 +1,22 @@ -// Test PGO when the kernel runs on a non-default device. -// The profile runtime should collect data from whichever device is active -// at exit time, which should be the device the user set. +// Test PGO when the kernel runs on a non-default device. The HSA-introspection +// drain walks every GPU agent at exit and collects device counters from each +// loaded code object, so it must still drain the device profile when the active +// device at exit is not device 0 (here the program selects device 1). // // REQUIRES: hip, amdgpu, multi-device +// Relies on the HSA agent-walk draining a device that is not current at exit; +// the Windows host-shadow drain only collects the current device. +// UNSUPPORTED: windows // RUN: %clang -x hip -fprofile-instr-generate -fcoverage-mapping \ // RUN: --offload-arch=%amdgpu_arch %s -o %t -L%hip_lib_path -lamdhip64 // RUN: env LLVM_PROFILE_FILE=%t.profraw \ // RUN: LD_LIBRARY_PATH=%hip_lib_path:$LD_LIBRARY_PATH \ // RUN: LLVM_PROFILE_VERBOSE=1 %run %t 2>&1 | FileCheck %s // -// The runtime should collect from the device that was set (device 1). -// CHECK: Collecting static profile data from device 1 -// CHECK: Successfully wrote device profile using shared API +// The drain must walk the agents and copy device counters even though the +// kernel ran on the non-default device. +// CHECK: Copied device sections: +// CHECK: walk complete: agents={{[0-9]+}} pairs={{[0-9]+}} found={{[0-9]+}} drained={{[1-9][0-9]*}} // CHECK: PASS #include diff --git a/compiler-rt/test/profile/device-pgo/README.md b/compiler-rt/test/profile/device-pgo/README.md new file mode 100644 index 0000000000000..61ee5153224fa --- /dev/null +++ b/compiler-rt/test/profile/device-pgo/README.md @@ -0,0 +1,105 @@ +# HIP device PGO / code coverage: standalone build & test recipe + +This directory provides a CMake-based recipe to build and exercise HIP device +profile-guided optimization (PGO) and source-based code coverage **outside +TheRock**, using only an `llvm-project` checkout plus a ROCm runtime. + +It builds, in one configure: + +- the host toolchain (`clang`, `clang++`, `lld`, `llvm-profdata`, `llvm-cov`) + and the lit-lite test utilities (`FileCheck`, `not`); +- the host ROCm drain runtime `clang_rt.profile_rocm` (opt-in, + `COMPILER_RT_BUILD_PROFILE_ROCM=ON`). On **Linux** this contains the HSA + introspection drain (`InstrProfilingPlatformROCm.cpp`); on **Windows** the + host-shadow drain (`InstrProfilingPlatformROCmWindows.cpp`); +- the **amdgcn device** profile runtime `libclang_rt.profile.a` (the baremetal + profile subset that provides `__llvm_profile_instrument_gpu` and the + `__llvm_profile_sections` bounds table), built for the `amdgcn-amd-amdhsa` + runtime target with LLVM libc for amdgcn. + +## Why a separate library + +Upstream relands HIP offload PGO runtime support as the **opt-in** +`clang_rt.profile_rocm` (llvm#201606), a `/MD` superset of `clang_rt.profile`; +the base library stays unchanged. The driver links `clang_rt.profile_rocm` +ahead of `clang_rt.profile` for HIP host links when profiling is requested +(see `clang/lib/Driver/ToolChains/{Linux,MSVC}.cpp`). This recipe just turns +the option on and builds the matching amdgcn device runtime. + +## Prerequisites + +- A ROCm installation (for `libamdhip64` and, on Linux, `libhsa-runtime64`), + e.g. `/opt/rocm`. Export `ROCM_PATH`. +- An AMD GPU visible to the runtime for the *run* step (the build step does + not need a GPU). `amdgpu-arch` should list your device(s). +- Ninja, a host C/C++ compiler, and Python 3. + +## Build + +```bash +export ROCM_PATH=/opt/rocm +./build.sh # builds into /build/device-pgo +# or: ./build.sh /path/to/builddir +``` + +Key outputs under the build dir: + +``` +bin/{clang,clang++,lld,llvm-profdata,llvm-cov,FileCheck,not} +lib/clang//lib//libclang_rt.profile_rocm.a +lib/clang//lib/amdgcn-amd-amdhsa/libclang_rt.profile.a +``` + +See `toolchain-cache.cmake` for the exact CMake variables, including the +`LLVM_RUNTIME_TARGETS="default;amdgcn-amd-amdhsa"` split. + +## Run the tests + +The lit-lite runner (`../run_gpu_tests.py`) compiles each `.hip` test with the +just-built toolchain, runs it on the GPU, and pipes output through `FileCheck`. +It auto-detects features (`multi-device` via `amdgpu-arch`) so tests that need +two visible GPUs are skipped on single-GPU hosts. + +```bash +python3 ../run_gpu_tests.py \ + --toolchain-bin "$PWD//bin" \ + --hip-lib-path "$ROCM_PATH/lib" \ + ../GPU ../AMDGPU +``` + +`--toolchain-bin` must be an **absolute** path (the runner executes each RUN +line from a temp directory). With the toolchain's `amdgpu-arch`/`offload-arch` +on hand, `--offload-arch=native` resolves automatically and the `multi-device` +feature is enabled when 2+ GPUs are visible (so multi-GPU tests run on a +multi-GPU host and are skipped otherwise). On a multi-gfx90a host this suite is +11 passed, 0 failed. + +## Manual workflow (for reference) + +```bash +CLANG=/bin/clang++ +# 1. Instrumented build (host + device). +$CLANG -O2 -fprofile-instr-generate -fcoverage-mapping \ + --offload-arch=gfx1100 -xhip app.hip -o app + +# 2. Run. Produces a host .profraw and a device +# .amdgcn-amd-amdhsa..profraw drained by clang_rt.profile_rocm. +LLVM_PROFILE_FILE='app-%p.profraw' ./app + +# 3. Merge (device profiles are merged per GPU arch). +/bin/llvm-profdata merge -o app.profdata app-*.profraw + +# 4. Coverage report (device). +/bin/llvm-cov show ./app -instr-profile=app.profdata +``` + +## Notes / environment-specific knobs + +- `--offload-arch` must match your GPU; the amdgcn device runtime is target + generic but the app's device code is per-arch. The build installs + `offload-arch` (and the `amdgpu-arch` alias) into `/bin`, so + `--offload-arch=native` works without a system ROCm `amdgpu-arch`. +- The amdgcn runtime target requires LLVM libc for amdgcn; if your environment + cannot build it, drop `libc` from + `RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES` only if your headers are + otherwise provided. diff --git a/compiler-rt/test/profile/device-pgo/build.sh b/compiler-rt/test/profile/device-pgo/build.sh new file mode 100755 index 0000000000000..edf90f42fb8c1 --- /dev/null +++ b/compiler-rt/test/profile/device-pgo/build.sh @@ -0,0 +1,56 @@ +#!/usr/bin/env bash +# Standalone (non-TheRock) build of the toolchain + host/device runtimes used by +# the HIP device-PGO / code-coverage tests. See toolchain-cache.cmake and +# README.md for details. +# +# ./build.sh [BUILD_DIR] +# +# Env knobs: +# LLVM_SRC path to the llvm-project checkout (default: repo root inferred +# from this script's location) +# JOBS parallelism for ninja (default: nproc) +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +# .../compiler-rt/test/profile/device-pgo -> repo root is four levels up. +LLVM_SRC="${LLVM_SRC:-$(cd "${SCRIPT_DIR}/../../../.." && pwd)}" +BUILD_DIR="${1:-${LLVM_SRC}/build/device-pgo}" +JOBS="${JOBS:-$(nproc)}" + +echo "llvm-project source : ${LLVM_SRC}" +echo "build directory : ${BUILD_DIR}" +echo "parallel jobs : ${JOBS}" + +cmake -G Ninja \ + -S "${LLVM_SRC}/llvm" \ + -B "${BUILD_DIR}" \ + -C "${SCRIPT_DIR}/toolchain-cache.cmake" + +# The 'clang' target also produces the clang++ symlink. The offload toolchain +# tools (clang-offload-bundler, clang-linker-wrapper, llvm-link, +# llvm-offload-binary) and offload-arch (also installed as amdgpu-arch) are +# needed to compile/link a HIP program and to resolve --offload-arch=native / +# the multi-device test feature. 'runtimes' builds both the host (default) and +# amdgcn device runtime targets. +ninja -C "${BUILD_DIR}" -j "${JOBS}" \ + clang lld \ + clang-offload-bundler clang-linker-wrapper llvm-link llvm-offload-binary \ + offload-arch \ + llvm-profdata llvm-cov FileCheck not \ + runtimes + +cat </lib//libclang_rt.profile_rocm.a +# lib/clang//lib/amdgcn-amd-amdhsa/libclang_rt.profile.a + +set(CMAKE_BUILD_TYPE Release CACHE STRING "") + +set(LLVM_ENABLE_PROJECTS "clang;lld" CACHE STRING "") +set(LLVM_ENABLE_RUNTIMES "compiler-rt" CACHE STRING "") +set(LLVM_TARGETS_TO_BUILD "host;AMDGPU" CACHE STRING "") +set(LLVM_ENABLE_PER_TARGET_RUNTIME_DIR ON CACHE BOOL "") +set(LLVM_ENABLE_ASSERTIONS ON CACHE BOOL "") + +set(CLANG_DEFAULT_LINKER "lld" CACHE STRING "") +set(CLANG_DEFAULT_RTLIB "compiler-rt" CACHE STRING "") + +# Make FileCheck / not available in the install/bin tree for the lit-lite runner. +set(LLVM_INSTALL_UTILS ON CACHE BOOL "") + +# Build host (default) and device (amdgcn) runtimes in one tree. +set(LLVM_RUNTIME_TARGETS "default;amdgcn-amd-amdhsa" CACHE STRING "") + +# Host runtimes: turn on the opt-in ROCm host drain library. It pulls in the +# sanitizer interception object libs, so sanitizers must be built too. +set(RUNTIMES_default_COMPILER_RT_BUILD_PROFILE_ROCM ON CACHE BOOL "") +set(RUNTIMES_default_COMPILER_RT_BUILD_SANITIZERS ON CACHE BOOL "") + +# Device runtime: the amdgcn baremetal profile subset, built with LLVM libc for +# amdgcn (freestanding C headers). +set(RUNTIMES_amdgcn-amd-amdhsa_CACHE_FILES + "${CMAKE_SOURCE_DIR}/../compiler-rt/cmake/caches/AMDGPU.cmake" CACHE STRING "") +set(RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES "compiler-rt;libc" CACHE STRING "") diff --git a/compiler-rt/test/profile/run_gpu_tests.py b/compiler-rt/test/profile/run_gpu_tests.py new file mode 100644 index 0000000000000..d339969ba245e --- /dev/null +++ b/compiler-rt/test/profile/run_gpu_tests.py @@ -0,0 +1,380 @@ +#!/usr/bin/env python3 +"""Minimal lit-style runner for the HIP device-PGO tests. + +The compiler-rt profile lit suite (and llvm-lit / FileCheck) is not part of the +installed ROCm artifact, but the toolchain, the amdgcn device profile runtime, +and the HIP runtime are. This runner executes the +``compiler-rt/test/profile/{GPU,AMDGPU}/*.hip`` tests directly against an +installed toolchain on a real GPU runner, interpreting just the slice of lit +markup those tests use: + + - ``// REQUIRES:`` / ``// UNSUPPORTED:`` boolean feature gating, + - ``// RUN:`` lines (with ``\\`` continuations) and the fixed substitution set + (%clang, %s, %t[.*], %amdgpu_arch, %hip_lib_path, %run, %%), + - delegation to ``FileCheck`` / ``not`` (real binaries if present on PATH, + otherwise shims backed by the ``filecheck`` PyPI package and a tiny + exit-code inverter). + +Each RUN line is executed via ``bash -e -o pipefail -c`` so pipes, redirection +and globbing behave as under lit. A test passes iff all its RUN lines exit 0. +""" + +import argparse +import os +import re +import shutil +import stat +import subprocess +import sys +import tempfile +from pathlib import Path + +# --- feature detection ------------------------------------------------------ + + +def _count_visible_gpus(toolchain_bin): + """Number of GPUs actually visible to the runtime, or 0 if unknown. + + Uses the toolchain's ``amdgpu-arch`` (one line per visible device). Unlike + the KFD topology under ``/sys/class/kfd`` this reflects what HIP/ROCr really + exposes -- it honours ``ROCR_VISIBLE_DEVICES`` / ``HIP_VISIBLE_DEVICES`` and + container device limits, so it matches what a test's ``hipGetDeviceCount`` + will see. It is also portable: Windows has no ``/dev/kfd``, but does ship + ``amdgpu-arch``. + """ + if not toolchain_bin: + return 0 + tb = Path(toolchain_bin) + exe = next( + (str(tb / c) for c in ("amdgpu-arch", "amdgpu-arch.exe") if (tb / c).exists()), + None, + ) + if exe is None: + return 0 + try: + proc = subprocess.run(exe, capture_output=True, text=True, timeout=60) + except (OSError, subprocess.SubprocessError): + return 0 + if proc.returncode != 0: + return 0 + return sum(1 for line in proc.stdout.splitlines() if line.strip()) + + +def detect_features(toolchain_bin=None, force=None): + """Return the set of lit features available on this runner. + + hip/amdgpu are assumed present (this runner only ever drives GPU tests on a + runner that has the toolchain + HIP). ``multi-device`` is derived from the + number of GPUs the runtime actually exposes (>= 2), via ``amdgpu-arch``. + """ + features = {"hip", "amdgpu"} + if sys.platform.startswith("linux"): + features.add("linux") + elif sys.platform.startswith("win"): + features.add("windows") + + if _count_visible_gpus(toolchain_bin) >= 2: + features.add("multi-device") + + if force: + for f in force: + features.add(f) + return features + + +# --- boolean expression evaluation (REQUIRES / UNSUPPORTED) ------------------ + +_TOKEN_RE = re.compile(r"\s*(\(|\)|\|\||&&|!|[\w.+-]+)\s*") + + +def _clause_to_py(clause): + out = [] + for tok in _TOKEN_RE.findall(clause): + if tok == "||": + out.append(" or ") + elif tok == "&&": + out.append(" and ") + elif tok == "!": + out.append(" not ") + elif tok in ("(", ")"): + out.append(tok) + elif tok == "true": + out.append("True") + elif tok == "false": + out.append("False") + else: + out.append("(%r in FEATURES)" % tok) + return "".join(out) or "True" + + +def eval_requires(expr, features): + """All comma-separated clauses must be true.""" + return all( + eval(_clause_to_py(c), {"__builtins__": {}}, {"FEATURES": features}) + for c in expr.split(",") + if c.strip() + ) + + +def eval_unsupported(expr, features): + """Unsupported if any comma-separated clause is true.""" + return any( + eval(_clause_to_py(c), {"__builtins__": {}}, {"FEATURES": features}) + for c in expr.split(",") + if c.strip() + ) + + +# --- test parsing ----------------------------------------------------------- + +_DIRECTIVE_RE = re.compile(r"(?://|#)\s*(RUN|REQUIRES|UNSUPPORTED):\s?(.*)") + + +def parse_test(path): + """Return (run_lines, requires, unsupported) for a test file.""" + runs, requires, unsupported = [], [], [] + cont = None + for raw in Path(path).read_text(errors="replace").splitlines(): + m = _DIRECTIVE_RE.search(raw) + if cont is not None: + # Continuation of a previous RUN line. + text = raw + cm = re.search(r"(?://|#)\s*RUN:\s?(.*)", raw) + if cm: + text = cm.group(1) + cont += " " + text.strip() + if cont.rstrip().endswith("\\"): + cont = cont.rstrip()[:-1] + else: + runs.append(cont) + cont = None + continue + if not m: + continue + kind, body = m.group(1), m.group(2) + if kind == "REQUIRES": + requires.append(body.strip()) + elif kind == "UNSUPPORTED": + unsupported.append(body.strip()) + elif kind == "RUN": + if body.rstrip().endswith("\\"): + cont = body.rstrip()[:-1] + else: + runs.append(body) + return runs, requires, unsupported + + +# --- substitutions ---------------------------------------------------------- + + +def make_substitutions(clang, clangxx, src, tprefix, arch, hip_lib_path): + # Order matters: longer / more specific tokens first; %% resolved last. + return [ + ("%clangxx", clangxx), + ("%clang", clang), + ("%amdgpu_arch", arch), + ("%hip_lib_path", hip_lib_path), + ("%run ", ""), + ("%s", str(src)), + ("%t", tprefix), + ("%%", "%"), + ] + + +def apply_substitutions(line, subs): + for token, value in subs: + line = line.replace(token, value) + return line + + +# --- tool shims (FileCheck / not) ------------------------------------------- + + +def ensure_tools(toolchain_bin, workdir): + """Build a PATH that resolves clang/llvm-*, FileCheck and not. + + Prefers real binaries under toolchain_bin; falls back to shims for FileCheck + (PyPI ``filecheck``) and ``not`` (exit-code inverter). + """ + shim_dir = workdir / "shims" + shim_dir.mkdir(parents=True, exist_ok=True) + path = os.pathsep.join([str(toolchain_bin), str(shim_dir), os.environ.get("PATH", "")]) + + def have(tool): + # File-based check (shutil.which is quirky across OSes / Git Bash). The + # shims are extensionless bash scripts, which Git Bash resolves via the + # shebang, so a real binary is anything matching tool or tool.exe. + tb = Path(toolchain_bin) + return (tb / tool).exists() or (tb / (tool + ".exe")).exists() + + def write_shim(name, body): + p = shim_dir / name + p.write_text(body) + p.chmod(p.stat().st_mode | stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH) + + if not have("FileCheck"): + write_shim( + "FileCheck", + '#!/usr/bin/env bash\n' + 'if command -v filecheck >/dev/null 2>&1; then exec filecheck "$@"; fi\n' + 'exec python3 -m filecheck "$@"\n', + ) + if not have("not"): + write_shim( + "not", + "#!/usr/bin/env bash\n" + 'if [ "$1" = "--crash" ]; then shift; "$@"; ec=$?; ' + "[ $ec -ge 128 ] && exit 0 || exit 1; fi\n" + '"$@"; ec=$?; [ $ec -eq 0 ] && exit 1 || exit 0\n', + ) + return path + + +# --- execution -------------------------------------------------------------- + + +def run_one(path, args, features, base_env): + runs, requires, unsupported = parse_test(path) + + for expr in requires: + if not eval_requires(expr, features): + return "UNSUPPORTED", "missing requirement: %s" % expr + for expr in unsupported: + if eval_unsupported(expr, features): + return "UNSUPPORTED", "unsupported: %s" % expr + if not runs: + return "UNSUPPORTED", "no RUN lines" + + workdir = Path(tempfile.mkdtemp(prefix="profgpu-")) + tprefix = str(workdir / "t") + subs = make_substitutions( + args.clang, args.clangxx, Path(path).resolve(), tprefix, args.amdgpu_arch, + args.hip_lib_path, + ) + + if args.dry_run: + print("# %s" % path) + for line in runs: + print(" " + apply_substitutions(line, subs).strip()) + return "DRYRUN", "" + + env = dict(base_env) + env["PATH"] = ensure_tools(Path(args.toolchain_bin), workdir) + timeout = args.timeout if args.timeout and args.timeout > 0 else None + for line in runs: + cmd = apply_substitutions(line, subs).strip() + try: + proc = subprocess.run( + ["bash", "-e", "-o", "pipefail", "-c", cmd], + cwd=str(workdir), + env=env, + capture_output=True, + text=True, + timeout=timeout, + ) + except subprocess.TimeoutExpired as e: + out = e.stdout or "" + err = e.stderr or "" + if isinstance(out, bytes): + out = out.decode("utf-8", "replace") + if isinstance(err, bytes): + err = err.decode("utf-8", "replace") + detail = "RUN timed out after %gs: %s\n%s%s" % ( + timeout, cmd, out, err, + ) + if not args.keep: + shutil.rmtree(workdir, ignore_errors=True) + return "FAIL", detail + if proc.returncode != 0: + detail = "RUN failed (rc=%d): %s\n%s%s" % ( + proc.returncode, cmd, proc.stdout, proc.stderr, + ) + if not args.keep: + shutil.rmtree(workdir, ignore_errors=True) + return "FAIL", detail + if not args.keep: + shutil.rmtree(workdir, ignore_errors=True) + return "PASS", "" + + +def discover(paths): + tests = [] + for p in paths: + p = Path(p) + if p.is_dir(): + tests.extend(sorted(str(x) for x in p.rglob("*.hip"))) + elif p.is_file(): + tests.append(str(p)) + return tests + + +def main(): + ap = argparse.ArgumentParser(description=__doc__) + ap.add_argument("tests", nargs="+", help="Test files or directories") + ap.add_argument("--toolchain-bin", required=False, + help="Directory with clang and llvm-* tools") + ap.add_argument("--hip-lib-path", default="", help="Directory with libamdhip64") + ap.add_argument("--amdgpu-arch", default="native") + ap.add_argument("--clang", help="Override clang path") + ap.add_argument("--clangxx", help="Override clang++ path") + ap.add_argument("--feature", action="append", default=[], + help="Force-enable an extra lit feature") + ap.add_argument("--dry-run", action="store_true", + help="Print resolved RUN lines without executing") + ap.add_argument("--keep", action="store_true", help="Keep per-test temp dirs") + ap.add_argument("--timeout", type=float, default=600, + help="Per-RUN-line timeout in seconds (<=0 disables); " + "guards against a hung GPU/compiler wedging the run") + args = ap.parse_args() + + if not args.dry_run and not args.toolchain_bin: + ap.error("--toolchain-bin is required unless --dry-run is given") + + if args.toolchain_bin: + binp = Path(args.toolchain_bin) + args.clang = args.clang or str(binp / "clang") + args.clangxx = args.clangxx or str(binp / "clang++") + else: + args.clang = args.clang or "clang" + args.clangxx = args.clangxx or "clang++" + + features = detect_features(args.toolchain_bin, args.feature) + print("# features: %s" % ", ".join(sorted(features))) + + base_env = dict(os.environ) + if args.toolchain_bin: + lib_dirs = [ + str(Path(args.toolchain_bin).parent / "lib"), # toolchain libs + ] + if args.hip_lib_path: + lib_dirs.append(args.hip_lib_path) + existing = base_env.get("LD_LIBRARY_PATH", "") + base_env["LD_LIBRARY_PATH"] = os.pathsep.join( + [d for d in lib_dirs if d] + ([existing] if existing else []) + ) + + tests = discover(args.tests) + if not tests: + print("error: no tests found", file=sys.stderr) + return 2 + + results = {"PASS": [], "FAIL": [], "UNSUPPORTED": [], "DRYRUN": []} + for t in tests: + status, detail = run_one(t, args, features, base_env) + results[status].append(t) + if status == "FAIL": + print("FAIL: %s" % t) + print(detail) + elif status in ("PASS", "UNSUPPORTED"): + print("%s: %s" % (status, t)) + + print( + "\n# summary: %d passed, %d failed, %d unsupported (of %d)" + % (len(results["PASS"]), len(results["FAIL"]), + len(results["UNSUPPORTED"]), len(tests)) + ) + return 1 if results["FAIL"] else 0 + + +if __name__ == "__main__": + sys.exit(main())