diff --git a/CMakeLists.txt b/CMakeLists.txt index 2d9c6aeace79..eb9fca039de9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -134,6 +134,33 @@ if (DEFINED ENV{ADDRESS_SANITIZER}) endif() ### End workaround +# Force every shared library we produce (rocMLIR's own and the embedded +# LLVM/MLIR `libLLVM*.so` / `libMLIR*.so`) to bind its internal function +# calls to its own definitions at link time. Intra-library calls then +# cannot be interposed by whichever libLLVM happens to be loaded first in +# the process. +# +# In practice this is defence-in-depth rather than a complete fix for the +# "two LLVMs, one process" problem. It removes PLT entries for all +# intra-library cl::* function calls, which reduces the attack surface, +# but it does NOT affect data-symbol interposition (vtables, RTTI, the +# process-global cl::SubCommand singleton pointer). Fully fixing the +# mlir-runner / xmir-runner JIT path against ROCm's `libLLVM.so.` +# still needs either a version script on `libLLVMSupport.so.git` +# or a repo-wide `CXX_VISIBILITY_PRESET=hidden` on the embedded LLVM; +# those are tracked separately. Applying -Bsymbolic-functions +# unconditionally is harmless (it never relaxes isolation, only +# tightens it) and forward-compatible with both of those follow-ups. +# +# The flag is a no-op on Windows (DLL imports already go through IAT) +# and unsupported on Apple's ld (skipped). +if (NOT WIN32 AND NOT APPLE) + set(CMAKE_SHARED_LINKER_FLAGS + "${CMAKE_SHARED_LINKER_FLAGS} -Wl,-Bsymbolic-functions") + set(CMAKE_MODULE_LINKER_FLAGS + "${CMAKE_MODULE_LINKER_FLAGS} -Wl,-Bsymbolic-functions") +endif() + # Set up the build for the LLVM/MLIR git-submodule include(cmake/llvm-project.cmake) diff --git a/cmake/llvm-project.cmake b/cmake/llvm-project.cmake index 816157b5b2b8..82a6b8f16d94 100644 --- a/cmake/llvm-project.cmake +++ b/cmake/llvm-project.cmake @@ -112,6 +112,20 @@ function(add_rocmlir_tool name) set(EXCLUDE_FROM_ALL ON) # LLVM functions read this variable, set it paranoidly endif() add_mlir_tool(${name} ${exclude_from_all} ${ARGN}) + + # Prevent symbols from static LLVM/MLIR archives linked into this tool + # from being re-exported to the dynamic symbol table. That matters most + # in static / fat-lib builds (BUILD_FAT_LIBROCKCOMPILER or BUILD_SHARED_LIBS + # OFF) where the tool pulls cl::opt definitions straight from libLLVMSupport.a + # -- without --exclude-libs,ALL the tool would unconditionally re-export + # them and any later-dlopened libLLVM.so.* (pulled in by libamdhip64 / + # libamd_comgr / runner libraries) would unify against them and trip + # "Option '...' already exists!" at static-init time. The flag is also + # a harmless no-op in the shared-lib build. Apple's ld does not accept + # GNU-style --exclude-libs. + if (NOT WIN32 AND NOT APPLE) + target_link_options(${name} PRIVATE "LINKER:--exclude-libs,ALL") + endif() endfunction() diff --git a/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmRuntimeLoader.h b/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmRuntimeLoader.h new file mode 100644 index 000000000000..98bdef66ed5c --- /dev/null +++ b/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmRuntimeLoader.h @@ -0,0 +1,130 @@ +//===- RocmRuntimeLoader.h - Lazy ROCm library loading utilities -*- C++-*-===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// Public API for delay-loading ROCm runtime shared libraries (libamdhip64, +// libhiprtc, libhsa-runtime64) from the MLIR ExecutionEngine and from +// downstream consumers. +// +// Consumers do NOT link the ROCm runtime at build time. Doing so would +// transitively pull `libamd_comgr` and ROCm's `libLLVM.so.` into +// the host process, which collides at static-init time with MLIR's own +// embedded LLVM (duplicate `cl::opt` registration aborts the process +// from `_dl_init` with "Option '...' already exists" or with a +// SmallPtrSet "Bucket < End" assertion). Loading these libraries with +// `dlmopen(LM_ID_NEWLM, ...)` on glibc puts them in a private link-map +// namespace where their LLVM cannot interpose ours. +// +// ROCm version compatibility: +// +// This loader is intentionally version-agnostic. It is built once +// and works against any ROCm major version present at runtime -- +// ROCm 4.x through any future ROCm release we have not yet seen. +// The selection algorithm prefers the unversioned SONAME (e.g. +// `libamdhip64.so` / `amdhip64.dll`), which is what every standard +// ROCm install ships and what `find_package(hip)` resolves; if that +// alias is absent (some runtime-only deployments), the loader falls +// back to enumerating versioned SONAMEs (`libamdhip64.so.` +// for descending MAJOR). HIP, HIPRTC and HSA each maintain a stable +// C ABI within a major version, so any HIP MAJOR the user has +// installed is acceptable to MLIR. There is no compile-time floor +// or ceiling on the ROCm version this code supports. +// +// Design choices that govern this API: +// +// - Header is platform-agnostic: no ``, no ``, no +// `_GNU_SOURCE` define. All platform-specific machinery lives in +// `RocmRuntimeLoader.cpp`. Downstream `add_mlir_library` users can +// include this header without inheriting Windows-macro pollution +// (`min`, `max`, `ERROR`, ...) or feature-test-macro surprises. +// +// - `LoadedLibrary` is an opaque struct rather than a `void *` typedef +// so a future change can carry extra state (search path used, debug +// info, ...) without breaking callers. +// +// - Cross-process coordination: `RocmSystemDetect` exports +// `mlirRocmSystemDetectGetHipHandle` (declared in +// `RocmSystemDetect.h`) so subsequent loaders share its HIP handle +// and the process keeps a single HSA session. KFD enforces one +// session per process; an independent second `dlmopen` would +// otherwise return `hipErrorNoDevice` from every call. +// +// - HIPRTC and HSA load into HIP's link-map namespace via the +// `relatedHandle` parameter so they share HIP's KFD session even +// when HIP itself was loaded into a non-default namespace. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_EXECUTIONENGINE_ROCMRUNTIMELOADER_H +#define MLIR_EXECUTIONENGINE_ROCMRUNTIMELOADER_H + +namespace mlir::rocm_loader { + +/// Identifies which ROCm shared library to delay-load. The enumerator +/// order is internal and may change; do not rely on it. +enum class Library { + Hip, + Hiprtc, + Hsa, +}; + +/// Opaque handle returned by `loadRocmLibrary`. `handle == nullptr` +/// indicates load failure; callers must treat that as "runtime +/// unavailable" and degrade gracefully. +struct LoadedLibrary { + void *handle = nullptr; +}; + +/// How `loadRocmLibrary` should coordinate with other loaders that +/// might already have opened the requested library in this process. +enum class CoordinationPolicy { + /// Default. For `Library::Hip`, attempt to reuse the HIP handle + /// owned by `RocmSystemDetect` (looked up via `RTLD_DEFAULT`); for + /// every other library this is equivalent to `Owned`. This is the + /// policy downstream consumers should use. + Auto, + + /// Skip the shared-handle lookup. The caller is the canonical + /// owner. Reserved for `RocmSystemDetect.cpp` to break recursion at + /// first load. + /// + /// IMPORTANT: do not use `Owned` from elsewhere. KFD permits only + /// one HSA session per process; on glibc each `Owned` call performs + /// a fresh `dlmopen(LM_ID_NEWLM, ...)` and thus opens HIP into a + /// new namespace. A second `Owned` invocation in the same process + /// will succeed at the `dlmopen` level but every subsequent HIP + /// call (`hipGetDeviceCount` etc.) returns `hipErrorNoDevice`. Use + /// `Auto` from non-canonical callers so they receive the shared + /// handle that `RocmSystemDetect` already holds. + Owned, +}; + +/// Load `lib` into a private link-map namespace and return an opaque +/// handle. On glibc this uses `dlmopen(LM_ID_NEWLM, ...)`; on other +/// POSIX platforms `dlopen(RTLD_LAZY | RTLD_LOCAL)`; on Windows +/// `LoadLibraryW` with UTF-8 -> UTF-16 conversion of the SONAME. +/// +/// When `relatedHandle` is non-null, the new library is opened in the +/// same link-map namespace as `relatedHandle` (glibc only; falls back +/// to the default namespace elsewhere). This is how HIPRTC and HSA +/// share HIP's KFD session. +/// +/// Returns a `LoadedLibrary` whose `handle` is null on failure. +/// Failures are non-fatal: this function never aborts the process. +LoadedLibrary +loadRocmLibrary(Library lib, void *relatedHandle = nullptr, + CoordinationPolicy policy = CoordinationPolicy::Auto); + +/// Resolve `name` in a previously-loaded library. Returns `nullptr` if +/// the library failed to load or if the symbol is absent. Callers +/// should treat `nullptr` as a soft error and disable the corresponding +/// feature. +void *resolveRocmSymbol(const LoadedLibrary &lib, const char *name); + +} // namespace mlir::rocm_loader + +#endif // MLIR_EXECUTIONENGINE_ROCMRUNTIMELOADER_H diff --git a/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmSystemDetect.h b/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmSystemDetect.h index 1d6a61c4ae21..714276f6ffba 100644 --- a/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmSystemDetect.h +++ b/external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmSystemDetect.h @@ -33,4 +33,37 @@ class RocmSystemDetect : public std::vector { } // namespace mlir +extern "C" { + +/// Returns the opaque HIP runtime handle owned by `RocmSystemDetect`, +/// or `nullptr` if HIP could not be loaded (or this binary did not +/// link `MLIRRocmExecutionEngineUtils`). +/// +/// `RocmSystemDetect` is the canonical owner of the per-process HIP +/// handle. When it loads `libamdhip64`, it uses +/// `dlmopen(LM_ID_NEWLM, ...)` (glibc) to put HIP and its transitive +/// dependencies (libamd_comgr, ROCm's libLLVM) in a private link-map +/// namespace. KFD enforces one HSA session per process; if a second +/// loader (for example mlir-runner's `libmlir_rocm_runtime.so`) opens +/// HIP into a *different* namespace, that second instance receives +/// `hipErrorNoDevice` from every call. To avoid that, all subsequent +/// HIP loaders look up this symbol via `RTLD_DEFAULT` and reuse the +/// returned handle. The recommended way to do that is to call +/// `mlir::rocm_loader::loadRocmLibrary(Library::Hip)` (defined in +/// `mlir/ExecutionEngine/RocmRuntimeLoader.h`), which performs the +/// lookup transparently. +/// +/// The function is `extern "C"` and uses an opaque `void *` so it can +/// be safely dlsym-ed from a TU that does not include this header +/// (notably from `RocmRuntimeLoader.cpp` itself, which avoids a +/// link-time dependency on `MLIRRocmExecutionEngineUtils`). +/// +/// Visibility: the symbol is published with `LLVM_ALWAYS_EXPORT` +/// (`__declspec(dllexport)` on Windows, default visibility on POSIX) +/// so it lands in the host process's dynamic symbol table for +/// `RTLD_DEFAULT` lookup. +void *mlirRocmSystemDetectGetHipHandle(); + +} // extern "C" + #endif // MLIR_EXECUTIONENGINE_ROCMSYSTEMDETECT_H_ diff --git a/external/llvm-project/mlir/lib/ExecutionEngine/CMakeLists.txt b/external/llvm-project/mlir/lib/ExecutionEngine/CMakeLists.txt index 79a82f493abc..1a04357d4192 100644 --- a/external/llvm-project/mlir/lib/ExecutionEngine/CMakeLists.txt +++ b/external/llvm-project/mlir/lib/ExecutionEngine/CMakeLists.txt @@ -12,6 +12,7 @@ set(LLVM_OPTIONAL_SOURCES ExecutionEngine.cpp Float16bits.cpp RocmDeviceName.cpp + RocmRuntimeLoader.cpp RocmRuntimeWrappers.cpp RunnerUtils.cpp OptUtils.cpp @@ -27,6 +28,38 @@ set(LLVM_OPTIONAL_SOURCES VulkanRuntime.h ) +# `MLIRRocmRuntimeLoader` consolidates the dlmopen / dlopen / +# LoadLibraryW scaffolding used by `mlir_rocm_runtime`, +# `MLIRRocmExecutionEngineUtils`, and downstream consumers (rocMLIR's +# `MLIRRockOps`, `rocmlir-tuning-driver`). +# +# `STATIC` is a hint that gets honored when `BUILD_SHARED_LIBS=OFF` +# (e.g. in the fat-`librockCompiler.a` build). In the default shared +# build, CMake produces `libMLIRRocmRuntimeLoader.so` and consumers +# pick it up as a `NEEDED` entry; that is harmless, because this +# library only depends on `LLVMSupport` (no ROCm runtime), so its +# transitive closure introduces no forbidden dependencies. +add_mlir_library(MLIRRocmRuntimeLoader + STATIC + RocmRuntimeLoader.cpp + + EXCLUDE_FROM_LIBMLIR + PARTIAL_SOURCES_INTENDED + + ADDITIONAL_HEADER_DIRS + ${MLIR_MAIN_INCLUDE_DIR}/mlir/ExecutionEngine + + LINK_COMPONENTS + Support +) +# Build with PIC so it can be linked into either shared libraries +# (`libmlir_rocm_runtime.so`) or executables (`rocmlir-tuning-driver`). +set_target_properties(MLIRRocmRuntimeLoader PROPERTIES + POSITION_INDEPENDENT_CODE ON) +# `${CMAKE_DL_LIBS}` is `-ldl` on Linux (for dlmopen / dlsym / +# dlinfo); empty on platforms that fold `dl` into libc. +target_link_libraries(MLIRRocmRuntimeLoader PUBLIC ${CMAKE_DL_LIBS}) + # Use a separate library for OptUtils, to avoid pulling in the entire JIT and # codegen infrastructure. Unlike MLIRExecutionEngine, this is part of # libMLIR.so. @@ -404,17 +437,20 @@ if(LLVM_ENABLE_PIC) endif() endif() + # Symbol suppression is now enforced at link time by the + # `mlir_rocm_runtime.map` version script applied below + # (`--version-script=...`). The dynsym is restricted to the `mgpu*` + # entry points; `llvm::EnableABIBreakingChecks` and other LLVM + # internals are not exported, so dlopen-time ODR collisions against + # a host-process LLVM cannot occur. `DISABLE_PCH_REUSE` is retained + # as a defence-in-depth measure for downstream consumers that + # disable the version script. add_mlir_library(mlir_rocm_runtime SHARED RocmRuntimeWrappers.cpp EXCLUDE_FROM_LIBMLIR - # TODO: this is merely a workaround. If this library depends on LLVMSupport, - # it should suppress symbols, or if it doesn't, it shouldn't link against - # it. This workaround prevents the library from defining the symbol - # llvm::EnableABIBreakingChecks, which would cause ODR-violations when - # dlopen-ed. DISABLE_PCH_REUSE ) @@ -447,51 +483,125 @@ if(LLVM_ENABLE_PIC) set_property(TARGET mlir_rocm_runtime PROPERTY INSTALL_RPATH_USE_LINK_PATH ON) - target_link_libraries(mlir_rocm_runtime - PUBLIC - hip::host hip::amdhip64 - ) - endif() + # `mlir_rocm_runtime` does NOT link `libamdhip64` at build time; it + # resolves HIP entry points via `dlmopen(LM_ID_NEWLM, ...)` on glibc + # (plain dlopen / LoadLibraryW elsewhere) so every transitive ROCm + # library -- most importantly `libamd_comgr` and its embedded + # `libLLVM.so.` -- lives in a private link-map namespace and + # cannot interpose the host process's LLVM at static-init time. + # + # We still need the HIP headers for POD struct layouts + # (`hipDeviceProp_t`, `hipMemcpyKind`, ...). `find_package(hip)` + # populates `hip_INCLUDE_DIR` which points at the HIP include tree; + # `hip::host` only carries compile-time definitions + # (`__HIP_PLATFORM_AMD__=1`). Wire both up without linking libamdhip64. + target_include_directories(mlir_rocm_runtime SYSTEM PRIVATE + ${hip_INCLUDE_DIR}) + target_compile_definitions(mlir_rocm_runtime PRIVATE + __HIP_PLATFORM_AMD__=1) + if (TARGET obj.mlir_rocm_runtime) + target_include_directories(obj.mlir_rocm_runtime SYSTEM PRIVATE + ${hip_INCLUDE_DIR}) + target_compile_definitions(obj.mlir_rocm_runtime PRIVATE + __HIP_PLATFORM_AMD__=1) + endif() + target_link_libraries(mlir_rocm_runtime PRIVATE MLIRRocmRuntimeLoader) + + # Strict symbol hiding. The wrapper only exposes a small C ABI named + # mgpu*. Anything else it picks up transitively (e.g. llvm::* helpers + # pulled in by inlined MLIR/LLVM headers, or EnableABIBreakingChecks) + # must stay internal. Otherwise the dynamic linker unifies those + # symbols with whichever LLVM instance wins the load order, which is + # exactly the source of the `cl::opt 'Option already exists'` crash + # when `rocmlir-driver`/`xmir-runner` also embeds LLVM. + # + # Implementation notes: `--version-script` overrides the global/local + # partitioning of already-visible symbols. We deliberately do NOT set + # CXX_VISIBILITY_PRESET=hidden because upstream + # `RocmRuntimeWrappers.cpp` does not decorate its mgpu* entry points + # with `visibility("default")`; a hidden preset would silently hide + # them at compile time and the version script could not re-expose + # them. Using default visibility + version-script gives us the + # correct export set (only mgpu*) without touching upstream source. + # `--exclude-libs,ALL` hardens against static archive symbols leaking + # into the dynsym in case this file is ever compiled into a static + # build. + if (NOT WIN32 AND NOT APPLE) + set(_mlir_rocm_runtime_vscript + "${CMAKE_CURRENT_SOURCE_DIR}/mlir_rocm_runtime.map") + target_link_options(mlir_rocm_runtime PRIVATE + "LINKER:--exclude-libs,ALL" + "LINKER:--version-script=${_mlir_rocm_runtime_vscript}") + set_property(TARGET mlir_rocm_runtime APPEND PROPERTY LINK_DEPENDS + "${_mlir_rocm_runtime_vscript}") + unset(_mlir_rocm_runtime_vscript) + endif() - add_mlir_library(MLIRRocmExecutionEngineUtils - SHARED - RocmSystemDetect.cpp + # `MLIRRocmExecutionEngineUtils` is the host-side ROCm helper used by + # `xmir-runner` to enumerate AMD devices via `RocmSystemDetect`. Its + # source `RocmSystemDetect.cpp` `#include`s ``, + # which means it needs `${hip_INCLUDE_DIR}` -- so it must live inside + # the `MLIR_ENABLE_ROCM_RUNNER` gate where `find_package(hip)` ran. + # In configurations with `MLIR_ENABLE_ROCM_RUNNER=OFF` (notably the + # `BUILD_FAT_LIBROCKCOMPILER=ON` build, which forces the runner off), + # this target is not built at all -- `xmir-runner` is also gated off + # in that mode, so nothing references it. The `MLIRRockUnitTests` + # CMake handles the missing target via `if (TARGET ...)`. + add_mlir_library(MLIRRocmExecutionEngineUtils + SHARED + RocmSystemDetect.cpp - ADDITIONAL_HEADER_DIRS - ${MLIR_MAIN_INCLUDE_DIR}/mlir/ExecutionEngine + ADDITIONAL_HEADER_DIRS + ${MLIR_MAIN_INCLUDE_DIR}/mlir/ExecutionEngine - DEPENDS - intrinsics_gen - ) + DEPENDS + intrinsics_gen + ) - if (CXX_SUPPORTS_CXX98_COMPAT_EXTRA_SEMI_FLAG) - target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE - "-Wno-c++98-compat-extra-semi") - endif() - if (CXX_SUPPORTS_WNO_RETURN_TYPE_C_LINKAGE_FLAG) - target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE - "-Wno-return-type-c-linkage") - endif() - if (CXX_SUPPORTS_WNO_NESTED_ANON_TYPES_FLAG) - target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE - "-Wno-nested-anon-types") - endif() - if (CXX_SUPPORTS_WNO_GNU_ANONYMOUS_STRUCT_FLAG) - target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE - "-Wno-gnu-anonymous-struct") - endif() + if (CXX_SUPPORTS_CXX98_COMPAT_EXTRA_SEMI_FLAG) + target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE + "-Wno-c++98-compat-extra-semi") + endif() + if (CXX_SUPPORTS_WNO_RETURN_TYPE_C_LINKAGE_FLAG) + target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE + "-Wno-return-type-c-linkage") + endif() + if (CXX_SUPPORTS_WNO_NESTED_ANON_TYPES_FLAG) + target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE + "-Wno-nested-anon-types") + endif() + if (CXX_SUPPORTS_WNO_GNU_ANONYMOUS_STRUCT_FLAG) + target_compile_options(obj.MLIRRocmExecutionEngineUtils PRIVATE + "-Wno-gnu-anonymous-struct") + endif() - target_link_libraries(MLIRRocmExecutionEngineUtils PRIVATE - MLIRExecutionEngineUtils - ) - # We do need both of these. One is for linking, to get the HIP - # functions. The other is for compiling, to get -D__HIP_PLATFORM_AMD__ - target_link_libraries(MLIRRocmExecutionEngineUtils PUBLIC - hip::host hip::amdhip64 - ) - target_link_libraries(obj.MLIRRocmExecutionEngineUtils PUBLIC - hip::host hip::amdhip64 - ) + target_link_libraries(MLIRRocmExecutionEngineUtils PRIVATE + MLIRExecutionEngineUtils + ) + + # MLIRRocmExecutionEngineUtils does NOT link `libamdhip64` at build + # time. Doing so would propagate the dependency to every consumer of + # this library (notably `xmir-runner`), which would then transitively + # pull in `libamd_comgr` and ROCm's `libLLVM.so.` and trigger + # the `cl::SubCommand` static-init collision against rocMLIR's + # embedded LLVM at process startup, before `main()` is even reached. + # `RocmSystemDetect.cpp` resolves the few HIP entry points it needs + # via `dlmopen(LM_ID_NEWLM, ...)` instead. We still need the HIP + # headers (for `hipDeviceProp_t`); wire those up via include dirs + # and `__HIP_PLATFORM_AMD__` only. + target_include_directories(MLIRRocmExecutionEngineUtils SYSTEM PRIVATE + ${hip_INCLUDE_DIR}) + target_compile_definitions(MLIRRocmExecutionEngineUtils PRIVATE + __HIP_PLATFORM_AMD__=1) + if (TARGET obj.MLIRRocmExecutionEngineUtils) + target_include_directories(obj.MLIRRocmExecutionEngineUtils SYSTEM PRIVATE + ${hip_INCLUDE_DIR}) + target_compile_definitions(obj.MLIRRocmExecutionEngineUtils PRIVATE + __HIP_PLATFORM_AMD__=1) + endif() + target_link_libraries(MLIRRocmExecutionEngineUtils PRIVATE + MLIRRocmRuntimeLoader) + endif() if(MLIR_ENABLE_SYCL_RUNNER OR MLIR_ENABLE_LEVELZERO_RUNNER) # Both runtimes require LevelZero, so we can find it once. diff --git a/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeLoader.cpp b/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeLoader.cpp new file mode 100644 index 000000000000..08fcd286ef2a --- /dev/null +++ b/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeLoader.cpp @@ -0,0 +1,255 @@ +//===- RocmRuntimeLoader.cpp - Lazy ROCm library loading utilities --------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// glibc's `dlmopen` is gated on `_GNU_SOURCE`. Define it before any system +// header is transitively included so the declaration is visible regardless of +// how the build picks compile flags. The upstream LLVM build sets `_GNU_SOURCE` +// repo-wide via `cmake/config-ix.cmake`, but downstream consumers compiling +// this TU through their own build system may not, so we define it defensively +// here. This define is confined to the implementation file and never leaks +// through the public header. +#if !defined(_WIN32) && !defined(_GNU_SOURCE) +#define _GNU_SOURCE +#endif + +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" + +#include "llvm/Support/ConvertUTF.h" +#include "llvm/Support/Debug.h" +#include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/raw_ostream.h" + +#include +#include +#include + +// We deliberately do NOT use `llvm::sys::DynamicLibrary` here even though it is +// the standard upstream wrapper for `dlopen`/`dlsym` and +// `LoadLibraryW`/`GetProcAddress`. The reason: on POSIX it always passes +// `RTLD_LAZY | RTLD_GLOBAL` to `dlopen` (see +// `lib/Support/Unix/DynamicLibrary.inc`), which is exactly the thing we need +// to avoid -- `RTLD_GLOBAL` lets the dynamic linker unify ROCm's `libLLVM.so` +// symbols with the host's embedded LLVM, which is the original `cl::opt` +// collision we are fixing. There is no public knob to swap in `RTLD_LOCAL` or +// `dlmopen(LM_ID_NEWLM, ...)`, and the namespace-isolation guarantee is the +// entire point of this loader. We therefore drop to the same OS APIs as +// `lib/Support/{Unix,Windows}/DynamicLibrary.inc` but with the flags we +// actually need. +#ifdef _WIN32 +#include +#else +#include +#endif + +#define DEBUG_TYPE "rocm-runtime-loader" + +namespace mlir::rocm_loader { +namespace { + +// Highest ROCm major version the loader will probe when iterating numeric +// SONAME suffixes (e.g. `libamdhip64.so.`). Picked generously so the loader +// continues working on future ROCm releases without code changes; bumping it +// has zero functional cost (a missing SONAME returns from `dlopen` in +// O(microseconds) on every modern libc, paid once at startup). The lower bound +// is `1` -- ROCm has never shipped a `.so.0`. Adjust upward if AMD ever reaches +// a major version above this constant. +constexpr unsigned kMaxProbedRocmMajor = 99; + +// Append `bare` (the unversioned alias, preferred when present), then +// `joiner(MAJOR)` for descending MAJOR, to `out`. The unversioned alias +// resolves through `LD_LIBRARY_PATH` / `RPATH` / `RUNPATH` / +// `/etc/ld.so.cache` on glibc (the user's expected policy, also what +// `find_package(hip)`, IREE and Triton do); the numeric fallback covers +// runtime-only installs where the symlink has been stripped. +// +// `joiner` produces the platform-decorated versioned name -- for HIP on POSIX +// it returns `libamdhip64.so.`, on Windows `amdhip64_.dll`, etc. +template +void appendCandidates(std::vector &out, llvm::StringRef bare, + Joiner joiner) { + out.emplace_back(bare.str()); + for (unsigned m = kMaxProbedRocmMajor; m >= 1; --m) + out.emplace_back(joiner(m)); +} + +// Build the SONAME candidate list for `lib`. On Windows, `Library::Hsa` +// returns an empty list because ROCm on Windows ships no HSA runtime; callers +// must treat `loadRocmLibrary(Hsa)` as "HSA unavailable" there. +// +// Windows HIPRTC has a quirk: AMD decorates the DLL name with the ROCm major +// AND minor (`hiprtc.dll`, e.g. `hiprtc0700.dll` for ROCm 7.0). For +// each candidate major we therefore probe `hiprtc00.dll`; AMD has only +// ever shipped the `.0` minor decoration in practice. Downstream consumers +// shipping a non-`.0` minor must put the DLL on `PATH` so the bare +// `hiprtc.dll` lookup picks it up. +std::vector candidatesFor(Library lib) { + std::vector out; + out.reserve(1 + kMaxProbedRocmMajor); + switch (lib) { + case Library::Hip: +#ifdef _WIN32 + appendCandidates(out, "amdhip64.dll", [](unsigned m) { + return "amdhip64_" + std::to_string(m) + ".dll"; + }); +#else + appendCandidates(out, "libamdhip64.so", [](unsigned m) { + return "libamdhip64.so." + std::to_string(m); + }); +#endif + return out; + case Library::Hiprtc: +#ifdef _WIN32 + appendCandidates(out, "hiprtc.dll", [](unsigned m) { + char buf[16]; + std::snprintf(buf, sizeof(buf), "hiprtc%02u00.dll", m); + return std::string(buf); + }); +#else + appendCandidates(out, "libhiprtc.so", [](unsigned m) { + return "libhiprtc.so." + std::to_string(m); + }); +#endif + return out; + case Library::Hsa: +#ifdef _WIN32 + return out; // empty: no HSA on Windows. +#else + // HSA's SONAME has been `.so.1` for the entire ROCm 4.x-7.x window, but + // we still iterate to be future-safe in case AMD ever bumps it. + appendCandidates(out, "libhsa-runtime64.so", [](unsigned m) { + return "libhsa-runtime64.so." + std::to_string(m); + }); + return out; +#endif + } + llvm_unreachable("unknown rocm_loader::Library enumerator"); +} + +// Open `path` so its symbols cannot interpose anything in the host process. +// On glibc this means `dlmopen(LM_ID_NEWLM, ...)` (a fresh link-map +// namespace); on other POSIX systems `dlopen(RTLD_LAZY | RTLD_LOCAL)`; on +// Windows `LoadLibraryW` (DLLs have private scopes per-DLL there). +// +// Returns null on failure -- never aborts. +void *openIsolated(const char *path) { +#ifdef _WIN32 + // Convert the UTF-8 SONAME to UTF-16 and call LoadLibraryW. This mirrors + // `llvm/lib/Support/Windows/DynamicLibrary.inc`. Our SONAMEs are pure ASCII + // today, but a downstream caller might extend the candidate list with a + // localized path, so we use the wide form unconditionally. + llvm::SmallVector wide; + if (!llvm::convertUTF8ToUTF16String(llvm::StringRef(path), wide)) { + LLVM_DEBUG(llvm::dbgs() + << DEBUG_TYPE ": bad UTF-8 in SONAME '" << path << "'\n"); + return nullptr; + } + HMODULE h = ::LoadLibraryW(reinterpret_cast(wide.data())); + if (!h) { + LLVM_DEBUG(llvm::dbgs() << DEBUG_TYPE ": LoadLibraryW(" << path + << ") failed (error " << ::GetLastError() << ")\n"); + } + return reinterpret_cast(h); +#else + // On glibc we open into a fresh link-map namespace so the loaded library's + // symbols cannot interpose the host's. Other POSIX libcs (musl, ...) lack + // `dlmopen`, so we settle for `RTLD_LOCAL`; isolation there is incomplete + // and depends on the host having hidden its own LLVM exports at link time + // (`-Wl,--exclude-libs,ALL`, visibility=hidden, ...). +#if defined(__GLIBC__) + void *h = ::dlmopen(LM_ID_NEWLM, path, RTLD_LAZY); +#else + void *h = ::dlopen(path, RTLD_LAZY | RTLD_LOCAL); +#endif + if (!h) { + LLVM_DEBUG(llvm::dbgs() << DEBUG_TYPE ": load failed for '" << path + << "': " << ::dlerror() << "\n"); + } + return h; +#endif +} + +// Open `path` into the SAME link-map namespace as `existingHandle`, so the +// new library shares state (most importantly KFD's per-process HSA session) +// with the previously-loaded HIP runtime. On glibc we look up +// `existingHandle`'s namespace via `dlinfo()` and pass it back to `dlmopen()`. +// On Windows / non-glibc POSIX, where namespaces don't exist, this is just a +// regular load. (`__GLIBC__` is never defined on Windows, so the single guard +// suffices.) +// +// Precondition: `existingHandle` is non-null. Caller routes the null case to +// `openIsolated`. +void *openInRelatedNamespace(const char *path, void *existingHandle) { +#if defined(__GLIBC__) + Lmid_t ns = LM_ID_NEWLM; + if (::dlinfo(existingHandle, RTLD_DI_LMID, &ns) != 0) + ns = LM_ID_NEWLM; + void *h = ::dlmopen(ns, path, RTLD_LAZY); + if (!h) { + LLVM_DEBUG(llvm::dbgs() << DEBUG_TYPE ": dlmopen(ns=" << ns << ", " << path + << ") failed: " << ::dlerror() << "\n"); + } + return h; +#else + (void)existingHandle; + return openIsolated(path); +#endif +} + +// Look up the HIP handle owned by `RocmSystemDetect`, if it has been loaded +// into this process. Returns null when the symbol is absent (typical for +// binaries that do not link `MLIRRocmExecutionEngineUtils`) or when +// `RocmSystemDetect` itself failed to load HIP. +// +// The lookup goes through `RTLD_DEFAULT` so we do not need a link-time +// dependency on `MLIRRocmExecutionEngineUtils`. On Windows, DLLs have private +// scopes so there is no equivalent coordination; ROCm-on-Windows also ships +// no HSA, so KFD's session limit does not apply. +void *getSharedHipHandle() { +#ifdef _WIN32 + return nullptr; +#else + using GetHandleFn = void *(*)(); + auto fn = reinterpret_cast( + ::dlsym(RTLD_DEFAULT, "mlirRocmSystemDetectGetHipHandle")); + return fn ? fn() : nullptr; +#endif +} + +} // namespace + +LoadedLibrary loadRocmLibrary(Library lib, void *relatedHandle, + CoordinationPolicy policy) { + LoadedLibrary out; + if (lib == Library::Hip && policy == CoordinationPolicy::Auto) { + if (void *shared = getSharedHipHandle()) { + out.handle = shared; + return out; + } + } + for (const std::string &cand : candidatesFor(lib)) { + out.handle = relatedHandle + ? openInRelatedNamespace(cand.c_str(), relatedHandle) + : openIsolated(cand.c_str()); + if (out.handle) + return out; + } + return out; +} + +void *resolveRocmSymbol(const LoadedLibrary &lib, const char *name) { + if (!lib.handle) + return nullptr; +#ifdef _WIN32 + return reinterpret_cast( + ::GetProcAddress(static_cast(lib.handle), name)); +#else + return ::dlsym(lib.handle, name); +#endif +} + +} // namespace mlir::rocm_loader diff --git a/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp b/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp index b984149ca6de..fe8bc76c2abe 100644 --- a/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp +++ b/external/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp @@ -10,16 +10,179 @@ // Also adds some debugging helpers that are helpful when writing MLIR code to // run on GPUs. // +// Linker discipline: this file does NOT link libamdhip64 at build time. +// The HIP runtime is loaded via the shared helpers in +// `mlir/ExecutionEngine/RocmRuntimeLoader.h`, which place libamdhip64 +// and its transitive dependencies (most importantly libamd_comgr and +// ROCm's libLLVM) into a private link-map namespace. See that header +// for the rationale (static-initializer collision between the two +// LLVMs, etc.). +// //===----------------------------------------------------------------------===// #include +#include +#include #include #include "mlir/ExecutionEngine/CRunnerUtils.h" +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" #include "llvm/ADT/ArrayRef.h" #include "hip/hip_runtime.h" +// Export tags match the upstream CudaRuntimeWrappers.cpp convention: every +// mgpu* entry point is default-visible so the JIT runner's dlsym() can +// find it. Internal helpers stay with file-default visibility (the CMake +// target applies a version script to enforce the export set). +#ifdef _WIN32 +#define MLIR_HIP_WRAPPERS_EXPORT __declspec(dllexport) +#else +#define MLIR_HIP_WRAPPERS_EXPORT __attribute__((visibility("default"))) +#endif + +namespace { + +namespace rocm_loader = ::mlir::rocm_loader; + +/// Table of resolved HIP entry points. All hipXXX wrappers further down +/// call through this table instead of linking against libamdhip64 at +/// build time. +struct HipSymbols { + rocm_loader::LoadedLibrary lib; + + const char *(*getErrorName)(hipError_t) = nullptr; + hipError_t (*moduleLoadData)(hipModule_t *, const void *) = nullptr; + hipError_t (*moduleUnload)(hipModule_t) = nullptr; + hipError_t (*moduleGetFunction)(hipFunction_t *, hipModule_t, + const char *) = nullptr; + hipError_t (*moduleLaunchKernel)(hipFunction_t, unsigned, unsigned, unsigned, + unsigned, unsigned, unsigned, unsigned, + hipStream_t, void **, void **) = nullptr; + hipError_t (*streamCreate)(hipStream_t *) = nullptr; + hipError_t (*streamDestroy)(hipStream_t) = nullptr; + hipError_t (*streamSynchronize)(hipStream_t) = nullptr; + hipError_t (*streamWaitEvent)(hipStream_t, hipEvent_t, unsigned) = nullptr; + hipError_t (*eventCreateWithFlags)(hipEvent_t *, unsigned) = nullptr; + hipError_t (*eventDestroy)(hipEvent_t) = nullptr; + hipError_t (*eventSynchronize)(hipEvent_t) = nullptr; + hipError_t (*eventRecord)(hipEvent_t, hipStream_t) = nullptr; + hipError_t (*malloc_)(void **, size_t) = nullptr; + hipError_t (*free_)(void *) = nullptr; + hipError_t (*memcpyAsync)(void *, const void *, size_t, hipMemcpyKind, + hipStream_t) = nullptr; + hipError_t (*memsetD32Async)(hipDeviceptr_t, int, size_t, + hipStream_t) = nullptr; + hipError_t (*memsetD16Async)(hipDeviceptr_t, short, size_t, + hipStream_t) = nullptr; + hipError_t (*hostRegister)(void *, size_t, unsigned) = nullptr; + hipError_t (*hostUnregister)(void *) = nullptr; + hipError_t (*hostGetDevicePointer)(void **, void *, unsigned) = nullptr; + hipError_t (*setDevice)(int) = nullptr; +}; + +HipSymbols loadHipSymbols() { + HipSymbols syms; + // `CoordinationPolicy::Auto` (the default) consults + // `mlirRocmSystemDetectGetHipHandle` first, so we share the process's + // single HSA session when RocmSystemDetect is present; otherwise we + // own the dlmopen ourselves. + syms.lib = rocm_loader::loadRocmLibrary(rocm_loader::Library::Hip); + if (!syms.lib.handle) { + std::fprintf( + stderr, "mlir_rocm_runtime: failed to load libamdhip64; hip calls will " + "fail. Ensure a ROCm install with libamdhip64.so is on " + "LD_LIBRARY_PATH / RPATH.\n"); + std::abort(); + } + +#define LOAD_HIP(FIELD, NAME, TYPE) \ + syms.FIELD = \ + reinterpret_cast(rocm_loader::resolveRocmSymbol(syms.lib, NAME)); \ + if (!syms.FIELD) { \ + std::fprintf(stderr, \ + "mlir_rocm_runtime: failed to resolve '%s' in " \ + "libamdhip64.\n", \ + NAME); \ + std::abort(); \ + } + + LOAD_HIP(getErrorName, "hipGetErrorName", const char *(*)(hipError_t)); + LOAD_HIP(moduleLoadData, "hipModuleLoadData", + hipError_t (*)(hipModule_t *, const void *)); + LOAD_HIP(moduleUnload, "hipModuleUnload", hipError_t (*)(hipModule_t)); + LOAD_HIP(moduleGetFunction, "hipModuleGetFunction", + hipError_t (*)(hipFunction_t *, hipModule_t, const char *)); + LOAD_HIP(moduleLaunchKernel, "hipModuleLaunchKernel", + hipError_t (*)(hipFunction_t, unsigned, unsigned, unsigned, unsigned, + unsigned, unsigned, unsigned, hipStream_t, void **, + void **)); + LOAD_HIP(streamCreate, "hipStreamCreate", hipError_t (*)(hipStream_t *)); + LOAD_HIP(streamDestroy, "hipStreamDestroy", hipError_t (*)(hipStream_t)); + LOAD_HIP(streamSynchronize, "hipStreamSynchronize", + hipError_t (*)(hipStream_t)); + LOAD_HIP(streamWaitEvent, "hipStreamWaitEvent", + hipError_t (*)(hipStream_t, hipEvent_t, unsigned)); + LOAD_HIP(eventCreateWithFlags, "hipEventCreateWithFlags", + hipError_t (*)(hipEvent_t *, unsigned)); + LOAD_HIP(eventDestroy, "hipEventDestroy", hipError_t (*)(hipEvent_t)); + LOAD_HIP(eventSynchronize, "hipEventSynchronize", hipError_t (*)(hipEvent_t)); + LOAD_HIP(eventRecord, "hipEventRecord", + hipError_t (*)(hipEvent_t, hipStream_t)); + LOAD_HIP(malloc_, "hipMalloc", hipError_t (*)(void **, size_t)); + LOAD_HIP(free_, "hipFree", hipError_t (*)(void *)); + LOAD_HIP( + memcpyAsync, "hipMemcpyAsync", + hipError_t (*)(void *, const void *, size_t, hipMemcpyKind, hipStream_t)); + LOAD_HIP(memsetD32Async, "hipMemsetD32Async", + hipError_t (*)(hipDeviceptr_t, int, size_t, hipStream_t)); + LOAD_HIP(memsetD16Async, "hipMemsetD16Async", + hipError_t (*)(hipDeviceptr_t, short, size_t, hipStream_t)); + LOAD_HIP(hostRegister, "hipHostRegister", + hipError_t (*)(void *, size_t, unsigned)); + LOAD_HIP(hostUnregister, "hipHostUnregister", hipError_t (*)(void *)); + LOAD_HIP(hostGetDevicePointer, "hipHostGetDevicePointer", + hipError_t (*)(void **, void *, unsigned)); + LOAD_HIP(setDevice, "hipSetDevice", hipError_t (*)(int)); + +#undef LOAD_HIP + return syms; +} + +const HipSymbols &getHip() { + static HipSymbols syms = loadHipSymbols(); + return syms; +} + +} // namespace + +// Redirect every bare hipXXX call-site below to go through the table. This +// keeps the rest of the file almost identical to upstream for easy merge. +#define hipGetErrorName(...) (::getHip().getErrorName(__VA_ARGS__)) +#define hipModuleLoadData(...) (::getHip().moduleLoadData(__VA_ARGS__)) +#define hipModuleUnload(...) (::getHip().moduleUnload(__VA_ARGS__)) +#define hipModuleGetFunction(...) (::getHip().moduleGetFunction(__VA_ARGS__)) +#define hipModuleLaunchKernel(...) (::getHip().moduleLaunchKernel(__VA_ARGS__)) +#define hipStreamCreate(...) (::getHip().streamCreate(__VA_ARGS__)) +#define hipStreamDestroy(...) (::getHip().streamDestroy(__VA_ARGS__)) +#define hipStreamSynchronize(...) (::getHip().streamSynchronize(__VA_ARGS__)) +#define hipStreamWaitEvent(...) (::getHip().streamWaitEvent(__VA_ARGS__)) +#define hipEventCreateWithFlags(...) \ + (::getHip().eventCreateWithFlags(__VA_ARGS__)) +#define hipEventDestroy(...) (::getHip().eventDestroy(__VA_ARGS__)) +#define hipEventSynchronize(...) (::getHip().eventSynchronize(__VA_ARGS__)) +#define hipEventRecord(...) (::getHip().eventRecord(__VA_ARGS__)) +#define hipMalloc(...) (::getHip().malloc_(__VA_ARGS__)) +#define hipFree(...) (::getHip().free_(__VA_ARGS__)) +#define hipMemcpyAsync(...) (::getHip().memcpyAsync(__VA_ARGS__)) +#define hipMemsetD32Async(...) (::getHip().memsetD32Async(__VA_ARGS__)) +#define hipMemsetD16Async(...) (::getHip().memsetD16Async(__VA_ARGS__)) +#define hipHostRegister(...) (::getHip().hostRegister(__VA_ARGS__)) +#define hipHostUnregister(...) (::getHip().hostUnregister(__VA_ARGS__)) +#define hipHostGetDevicePointer(...) \ + (::getHip().hostGetDevicePointer(__VA_ARGS__)) +#define hipSetDevice(...) (::getHip().setDevice(__VA_ARGS__)) + #define HIP_REPORT_IF_ERROR(expr) \ [](hipError_t result) { \ if (!result) \ @@ -32,23 +195,25 @@ thread_local static int32_t defaultDevice = 0; -extern "C" hipModule_t mgpuModuleLoad(void *data, size_t /*gpuBlobSize*/) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT hipModule_t +mgpuModuleLoad(void *data, size_t /*gpuBlobSize*/) { hipModule_t module = nullptr; HIP_REPORT_IF_ERROR(hipModuleLoadData(&module, data)); return module; } -extern "C" hipModule_t mgpuModuleLoadJIT(void *data, int optLevel) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT hipModule_t +mgpuModuleLoadJIT(void *data, int optLevel) { assert(false && "This function is not available in HIP."); return nullptr; } -extern "C" void mgpuModuleUnload(hipModule_t module) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuModuleUnload(hipModule_t module) { HIP_REPORT_IF_ERROR(hipModuleUnload(module)); } -extern "C" hipFunction_t mgpuModuleGetFunction(hipModule_t module, - const char *name) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT hipFunction_t +mgpuModuleGetFunction(hipModule_t module, const char *name) { hipFunction_t function = nullptr; HIP_REPORT_IF_ERROR(hipModuleGetFunction(&function, module, name)); return function; @@ -57,78 +222,83 @@ extern "C" hipFunction_t mgpuModuleGetFunction(hipModule_t module, // The wrapper uses intptr_t instead of ROCM's unsigned int to match // the type of MLIR's index type. This avoids the need for casts in the // generated MLIR code. -extern "C" void mgpuLaunchKernel(hipFunction_t function, intptr_t gridX, - intptr_t gridY, intptr_t gridZ, - intptr_t blockX, intptr_t blockY, - intptr_t blockZ, int32_t smem, - hipStream_t stream, void **params, - void **extra, size_t /*paramsCount*/) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuLaunchKernel(hipFunction_t function, intptr_t gridX, intptr_t gridY, + intptr_t gridZ, intptr_t blockX, intptr_t blockY, + intptr_t blockZ, int32_t smem, hipStream_t stream, + void **params, void **extra, size_t /*paramsCount*/) { HIP_REPORT_IF_ERROR(hipModuleLaunchKernel(function, gridX, gridY, gridZ, blockX, blockY, blockZ, smem, stream, params, extra)); } -extern "C" hipStream_t mgpuStreamCreate() { +extern "C" MLIR_HIP_WRAPPERS_EXPORT hipStream_t mgpuStreamCreate() { hipStream_t stream = nullptr; HIP_REPORT_IF_ERROR(hipStreamCreate(&stream)); return stream; } -extern "C" void mgpuStreamDestroy(hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuStreamDestroy(hipStream_t stream) { HIP_REPORT_IF_ERROR(hipStreamDestroy(stream)); } -extern "C" void mgpuStreamSynchronize(hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuStreamSynchronize(hipStream_t stream) { return HIP_REPORT_IF_ERROR(hipStreamSynchronize(stream)); } -extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuStreamWaitEvent(hipStream_t stream, + hipEvent_t event) { HIP_REPORT_IF_ERROR(hipStreamWaitEvent(stream, event, /*flags=*/0)); } -extern "C" hipEvent_t mgpuEventCreate() { +extern "C" MLIR_HIP_WRAPPERS_EXPORT hipEvent_t mgpuEventCreate() { hipEvent_t event = nullptr; HIP_REPORT_IF_ERROR(hipEventCreateWithFlags(&event, hipEventDisableTiming)); return event; } -extern "C" void mgpuEventDestroy(hipEvent_t event) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuEventDestroy(hipEvent_t event) { HIP_REPORT_IF_ERROR(hipEventDestroy(event)); } -extern "C" void mgpuEventSynchronize(hipEvent_t event) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuEventSynchronize(hipEvent_t event) { HIP_REPORT_IF_ERROR(hipEventSynchronize(event)); } -extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuEventRecord(hipEvent_t event, + hipStream_t stream) { HIP_REPORT_IF_ERROR(hipEventRecord(event, stream)); } -extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/, - bool /*isHostShared*/) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void *mgpuMemAlloc(uint64_t sizeBytes, + hipStream_t /*stream*/, + bool /*isHostShared*/) { void *ptr; HIP_REPORT_IF_ERROR(hipMalloc(&ptr, sizeBytes)); return ptr; } -extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuMemFree(void *ptr, + hipStream_t /*stream*/) { HIP_REPORT_IF_ERROR(hipFree(ptr)); } -extern "C" void mgpuMemcpy(void *dst, void *src, size_t sizeBytes, - hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuMemcpy(void *dst, void *src, size_t sizeBytes, hipStream_t stream) { HIP_REPORT_IF_ERROR( hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); } -extern "C" void mgpuMemset32(void *dst, int value, size_t count, - hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuMemset32(void *dst, int value, size_t count, hipStream_t stream) { HIP_REPORT_IF_ERROR(hipMemsetD32Async(reinterpret_cast(dst), value, count, stream)); } -extern "C" void mgpuMemset16(void *dst, int short value, size_t count, - hipStream_t stream) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuMemset16(void *dst, int short value, size_t count, hipStream_t stream) { HIP_REPORT_IF_ERROR(hipMemsetD16Async(reinterpret_cast(dst), value, count, stream)); } @@ -137,13 +307,14 @@ extern "C" void mgpuMemset16(void *dst, int short value, size_t count, // Allows to register byte array with the ROCM runtime. Helpful until we have // transfer functions implemented. -extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void +mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0)); } // Allows to register a MemRef with the ROCm runtime. Helpful until we have // transfer functions implemented. -extern "C" void +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType *descriptor, int64_t elementSizeBytes) { @@ -165,15 +336,15 @@ mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType *descriptor, mgpuMemHostRegister(ptr, sizeBytes); } -// Allows to unregister byte array with the ROCM runtime. Helpful until we have +// Allows to unregister byte array with the ROCm runtime. Helpful until we have // transfer functions implemented. -extern "C" void mgpuMemHostUnregister(void *ptr) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuMemHostUnregister(void *ptr) { HIP_REPORT_IF_ERROR(hipHostUnregister(ptr)); } // Allows to unregister a MemRef with the ROCm runtime. Helpful until we have // transfer functions implemented. -extern "C" void +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuMemHostUnregisterMemRef(int64_t rank, StridedMemRefType *descriptor, int64_t elementSizeBytes) { @@ -188,7 +359,7 @@ void mgpuMemGetDevicePointer(T *hostPtr, T **devicePtr) { hipHostGetDevicePointer((void **)devicePtr, hostPtr, /*flags=*/0)); } -extern "C" StridedMemRefType +extern "C" MLIR_HIP_WRAPPERS_EXPORT StridedMemRefType mgpuMemGetDeviceMemRef1dFloat(float *allocated, float *aligned, int64_t offset, int64_t size, int64_t stride) { float *devicePtr = nullptr; @@ -196,7 +367,7 @@ mgpuMemGetDeviceMemRef1dFloat(float *allocated, float *aligned, int64_t offset, return {devicePtr, devicePtr, offset, {size}, {stride}}; } -extern "C" StridedMemRefType +extern "C" MLIR_HIP_WRAPPERS_EXPORT StridedMemRefType mgpuMemGetDeviceMemRef1dInt32(int32_t *allocated, int32_t *aligned, int64_t offset, int64_t size, int64_t stride) { int32_t *devicePtr = nullptr; @@ -204,7 +375,7 @@ mgpuMemGetDeviceMemRef1dInt32(int32_t *allocated, int32_t *aligned, return {devicePtr, devicePtr, offset, {size}, {stride}}; } -extern "C" void mgpuSetDefaultDevice(int32_t device) { +extern "C" MLIR_HIP_WRAPPERS_EXPORT void mgpuSetDefaultDevice(int32_t device) { defaultDevice = device; HIP_REPORT_IF_ERROR(hipSetDevice(device)); } diff --git a/external/llvm-project/mlir/lib/ExecutionEngine/RocmSystemDetect.cpp b/external/llvm-project/mlir/lib/ExecutionEngine/RocmSystemDetect.cpp index 0e0c539d21f3..c19d1365ba09 100644 --- a/external/llvm-project/mlir/lib/ExecutionEngine/RocmSystemDetect.cpp +++ b/external/llvm-project/mlir/lib/ExecutionEngine/RocmSystemDetect.cpp @@ -6,31 +6,118 @@ // //===----------------------------------------------------------------------===// // -// This file implements the system detection of ROCm devices on the current -// system. +// Detects ROCm devices on the current system without link-time +// dependencies on `libamdhip64`. The HIP entry points are resolved via +// the shared helpers in `mlir/ExecutionEngine/RocmRuntimeLoader.h`, +// which hide the runtime in a private link-map namespace and therefore +// keep ROCm's libLLVM out of the host process's LLVM scope. +// +// This translation unit is also the canonical *owner* of the HIP +// handle: it exposes the handle through the extern-C function +// `mlirRocmSystemDetectGetHipHandle()` so other loaders +// (`libmlir_rocm_runtime.so`, rocMLIR's `MLIRRockOps`, the tuning +// driver) can reuse it. KFD only allows one user-space HSA session per +// process; a second `dlmopen(LM_ID_NEWLM, ...)` would end up in a +// different namespace and every call from it would return +// `hipErrorNoDevice`. // //===----------------------------------------------------------------------===// #include "mlir/ExecutionEngine/RocmSystemDetect.h" #include "mlir/ExecutionEngine/RocmDeviceName.h" +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" -#include "llvm/Support/Error.h" +#include "llvm/Support/Compiler.h" +#if defined(__GNUC__) && !defined(__clang__) #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wpedantic" +#endif #include "hip/hip_runtime.h" +#if defined(__GNUC__) && !defined(__clang__) #pragma GCC diagnostic pop +#endif -#define DEBUG_TYPE "execution-engine-rocm-system-detect" +#include using namespace mlir; #define TO_STR(x) llvm::StringRef(std::to_string(x)) +namespace { + +/// Function-pointer table for the (very small) HIP surface this file +/// uses. A null `handle` means HIP could not be loaded at all; the +/// constructor then treats the system as having zero GPUs. +struct HipSymbols { + rocm_loader::LoadedLibrary lib; + hipError_t (*getDeviceCount)(int *) = nullptr; + hipError_t (*getDeviceProperties)(hipDeviceProp_t *, int) = nullptr; +}; + +const HipSymbols &getHip() { + static HipSymbols syms = []() { + HipSymbols s; + // `CoordinationPolicy::Owned`: RocmSystemDetect is the canonical + // owner of the HIP handle. We must not consult + // `mlirRocmSystemDetectGetHipHandle` here; that symbol points back + // at *us* and would be uninitialised at this point. + s.lib = rocm_loader::loadRocmLibrary( + rocm_loader::Library::Hip, /*relatedHandle=*/nullptr, + rocm_loader::CoordinationPolicy::Owned); + if (!s.lib.handle) { + std::fprintf(stderr, + "RocmSystemDetect: libamdhip64 not found on the loader " + "search path; ROCm device detection disabled.\n"); + return s; + } + s.getDeviceCount = reinterpret_cast( + rocm_loader::resolveRocmSymbol(s.lib, "hipGetDeviceCount")); + // ROCm 6+ ABI-stable variant; fall back to the legacy symbol on + // older installs. + s.getDeviceProperties = + reinterpret_cast( + rocm_loader::resolveRocmSymbol(s.lib, + "hipGetDevicePropertiesR0600")); + if (!s.getDeviceProperties) { + s.getDeviceProperties = + reinterpret_cast( + rocm_loader::resolveRocmSymbol(s.lib, "hipGetDeviceProperties")); + } + if (!s.getDeviceCount || !s.getDeviceProperties) { + std::fprintf(stderr, + "RocmSystemDetect: libamdhip64 loaded but required " + "symbols are missing; ROCm device detection disabled.\n"); + s.lib.handle = nullptr; + } + return s; + }(); + return syms; +} + +} // namespace + +// Cross-library coordination export. The full contract lives on the +// declaration in `RocmSystemDetect.h`; `LLVM_ALWAYS_EXPORT` publishes +// the symbol in the host process's dynamic symbol table so other +// loaders can find it via `RTLD_DEFAULT` (POSIX) or its Windows +// equivalent. The macro expands to `__declspec(dllexport)` on +// Windows, `[[gnu::visibility("default")]]` / +// `__attribute__((visibility("default")))` on POSIX, and is the +// standard upstream way to mark a symbol as forcibly external; see +// `llvm/Support/Compiler.h` for the full definition. +extern "C" LLVM_ALWAYS_EXPORT void *mlirRocmSystemDetectGetHipHandle() { + return getHip().lib.handle; +} + RocmSystemDetect::RocmSystemDetect() { + const HipSymbols &hip = getHip(); + if (!hip.lib.handle) + return; + // collect all GPUs int count = 0; - hipError_t herr = hipGetDeviceCount(&count); + hipError_t herr = hip.getDeviceCount(&count); if (herr != hipSuccess) { llvm::errs() << "hipGetDeviceCount() should never fail\n"; return; @@ -38,7 +125,7 @@ RocmSystemDetect::RocmSystemDetect() { for (int i = 0; i < count; ++i) { hipDeviceProp_t deviceProps; - herr = hipGetDeviceProperties(&deviceProps, i); + herr = hip.getDeviceProperties(&deviceProps, i); if (herr == hipSuccess) { RocmDeviceName arch; if (succeeded(arch.parse(deviceProps.gcnArchName))) { diff --git a/external/llvm-project/mlir/lib/ExecutionEngine/mlir_rocm_runtime.map b/external/llvm-project/mlir/lib/ExecutionEngine/mlir_rocm_runtime.map new file mode 100644 index 000000000000..524a08c80ea7 --- /dev/null +++ b/external/llvm-project/mlir/lib/ExecutionEngine/mlir_rocm_runtime.map @@ -0,0 +1,13 @@ +# Linker version script for libmlir_rocm_runtime.so. +# +# The runtime wrapper exposes a small C ABI named with the "mgpu" prefix. +# Hide everything else (in particular, any llvm::* / MLIR:: symbols that a +# non-trivial link accidentally pulls in via hip/clang-rt objects) so the +# dynamic linker cannot unify them with rocMLIR's embedded LLVM or with +# ROCm's system libLLVM.so when both end up in the same address space. +{ + global: + mgpu*; + local: + *; +}; diff --git a/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h b/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h index 8835dd96235e..2b23e0f84a12 100644 --- a/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h +++ b/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h @@ -16,6 +16,8 @@ #include "mlir/Support/LLVM.h" #include "llvm/ADT/ArrayRef.h" +#include + namespace mlir { namespace rock { /// A structure containing information about a given AMD chip's features @@ -117,6 +119,18 @@ AmdArchInfo lookupArchInfo(StringRef arch); bool isDirectToLDSSupported(GemmFeatures features); bool isGlobalPrefetchSupported(StringRef arch); bool isAsyncDirectToLDSSupported(StringRef arch); + +/// Number of AMD GPUs visible via the HIP runtime, or 0 if HIP cannot be +/// loaded or returns an error. `MLIRRockOps` does not link `libamdhip64` at +/// build time; this query delay-loads it via `dlopen` / `LoadLibraryW` +/// internally. Callers must be prepared for the runtime to be missing (for +/// example in CI containers that have no ROCm installed). +unsigned nativeDeviceCount(); + +/// Hardware-reported `gcnArchName` for the given device id (e.g. "gfx942"), +/// or the empty string if the runtime is unavailable / the device id is +/// invalid. Loads HIP lazily, just like `nativeDeviceCount()`. +std::string nativeArchName(unsigned deviceId); } // namespace rock } // namespace mlir diff --git a/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp b/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp index decd04bc3209..9ab0ee2ed1b3 100644 --- a/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp +++ b/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp @@ -1,10 +1,23 @@ -//===- AmdArchDb.cpp - Dtabase of AMD GPU features ------------------===// +//===- AmdArchDb.cpp - Database of AMD GPU features -----------------------===// // // Part of the MLIR 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 // //===----------------------------------------------------------------------===// +// +// `MLIRRockOps` does not link `libamdhip64` / `libhsa-runtime64` at build +// time. Doing so would transitively pull in ROCm's `libamd_comgr` and its +// embedded `libLLVM.so.*`, which collides with rocMLIR's own LLVM at load +// time (duplicate `cl::opt` registration, corrupted global command-line +// parser state, ...). The HIP and HSA symbols used by +// `rock.arch = "native[:N]"` are resolved on demand via the shared +// `mlir::rocm_loader` helpers; see +// `external/llvm-project/mlir/include/mlir/ExecutionEngine/RocmRuntimeLoader.h` +// for the full rationale, including how we share a single HSA session +// across all consumers of libamdhip64 in the process. +// +//===----------------------------------------------------------------------===// #include "mlir/Dialect/Rock/IR/AmdArchDb.h" @@ -12,18 +25,28 @@ #include "mlir/Dialect/Rock/IR/RockGemmGemmWrapperInterface.h" #include "mlir/Dialect/Rock/IR/RockGemmWrapperInterface.h" #include "mlir/Dialect/Rock/IR/RockTypes.h" +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/TypeUtilities.h" #include "llvm/ADT/ArrayRef.h" - +#include "llvm/ADT/DenseMap.h" #include "llvm/ADT/StringSwitch.h" #include "llvm/MC/TargetRegistry.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" -// HIP and HSA are not supported on Windows CI. -#ifndef _WIN32 +#include +#include +#include + +// HIP / HSA headers are pulled in for their POD types (`hipDeviceProp_t`, +// `hsa_agent_t`, ...); the shared libraries themselves are loaded at +// runtime by `mlir::rocm_loader`, not linked at build time. +// `__HIP_PLATFORM_AMD__` picks the AMD variant of `hipDeviceProp_t`. +#define __HIP_PLATFORM_AMD__ 1 #include "hip/hip_runtime_api.h" + +#ifndef _WIN32 #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" #endif @@ -129,6 +152,17 @@ static constexpr AmdArchInfo /*hasOcpFp8ConversionInstrs=*/true, /*hasScaledGemm=*/false, /*maxNumXCC=*/1, /*hasLdsTransposeLoad=*/false); +// Parse one of the rock-arch syntaxes into `(chip, deviceId)`: +// - "native" -> ("native", 0) +// - "native:" -> ("native", N) for a non-negative integer N that +// fits in `unsigned` +// - "gfx<...>" -> ("gfx<...>", 0) +// - ":gfx<...>" -> ("gfx<...>", 0) +// +// Anything malformed -- in particular `"native:foo"`, `"native:1abc"`, +// `"native:"` -- aborts the process with a fatal error. Silently treating +// such input as `"native:0"` would mask user error and target the wrong +// GPU on multi-GPU systems. static std::tuple parseArchString(StringRef arch) { std::tuple ret("", 0); @@ -136,9 +170,23 @@ static std::tuple parseArchString(StringRef arch) { std::tie(firstPart, remainingParts) = arch.split(':'); if (firstPart == "native") { std::get<0>(ret) = firstPart; - if (unsigned long long deviceId; - !llvm::getAsUnsignedInteger(remainingParts, 0, deviceId)) { - std::get<1>(ret) = deviceId; + // `StringRef::split(':')` returns `("native", "")` for BOTH `"native"` + // (no separator) and `"native:"` (separator at end). To tell them apart + // we re-check the original input for the presence of a colon: if a colon + // was present, the suffix is mandatory (and must parse), otherwise the + // implicit deviceId is 0. + bool hasSeparator = arch.contains(':'); + if (hasSeparator) { + unsigned long long deviceId = 0; + if (remainingParts.empty() || + llvm::getAsUnsignedInteger(remainingParts, 0, deviceId) || + deviceId > std::numeric_limits::max()) + llvm::report_fatal_error( + llvm::Twine("Invalid `rock.arch = \"native:") + remainingParts + + "\"`: the suffix after `native:` must be a non-negative integer " + "device id (got `" + + remainingParts + "`)."); + std::get<1>(ret) = static_cast(deviceId); } } else { auto chipPos = firstPart.find("gfx"); @@ -153,216 +201,340 @@ static std::tuple parseArchString(StringRef arch) { return ret; } -// native arch is not supported in Windows, which lacks both HSA and -// HIP libraries during CI. For more information check: -// https://github.com/ROCm/rocMLIR/pull/1790 -#ifndef _WIN32 namespace { -template -std::enable_if_t, void> -checkAndSetInfo(StringRef name, LHS &lhs, RHS &&rhs) { - if (lhs != static_cast(rhs)) { - LLVM_DEBUG(llvm::dbgs() << "NOTE: Value discrepancy for " << name << ": " - << lhs << " (old) != " << rhs - << " (new). Proceeding with " << rhs << ".\n"); - lhs = std::forward(rhs); +//===----------------------------------------------------------------------===// +// HIP delay-load. +//===----------------------------------------------------------------------===// + +/// Function-pointer table resolved from `libamdhip64`. A null `handle` +/// means HIP was not loaded; callers must check before use. The handle +/// is intentionally leaked at process teardown: anything HIP pulled in +/// (most notably `libamd_comgr` and ROCm's `libLLVM.so`) must stay +/// mapped while any pointer HIP returned is still in flight. +struct HipRuntime { + rocm_loader::LoadedLibrary lib; + hipError_t (*getDeviceCount)(int *) = nullptr; + hipError_t (*getDeviceProperties)(hipDeviceProp_t *, int) = nullptr; +}; + +HipRuntime loadHipRuntime() { + HipRuntime rt; + rt.lib = rocm_loader::loadRocmLibrary(rocm_loader::Library::Hip); + if (!rt.lib.handle) { + LLVM_DEBUG(llvm::dbgs() + << "rock-amd-arch-db: libamdhip64 not found on the loader " + "search path; disabling native-arch detection\n"); + return rt; } + + rt.getDeviceCount = reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, "hipGetDeviceCount")); + // HIP 6+ renamed the struct-stable form to `...R0600`. Prefer that if + // available; fall back to the legacy symbol for older ROCm installs. + rt.getDeviceProperties = + reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, + "hipGetDevicePropertiesR0600")); + if (!rt.getDeviceProperties) + rt.getDeviceProperties = + reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, "hipGetDeviceProperties")); + + if (!rt.getDeviceCount || !rt.getDeviceProperties) { + LLVM_DEBUG(llvm::dbgs() + << "rock-amd-arch-db: HIP loaded but required symbols are " + "missing; disabling native-arch detection\n"); + rt.lib.handle = nullptr; + } + return rt; +} + +const HipRuntime &getHipRuntime() { + static HipRuntime rt = loadHipRuntime(); + return rt; } -struct AgentInfo { - // Input fields: - // The ID of the GPU device that we are looking for. - unsigned deviceId; - // Used in acquireAgentInfo, to compute GPU internal IDs. +//===----------------------------------------------------------------------===// +// HSA delay-load. +//===----------------------------------------------------------------------===// + +#ifndef _WIN32 +/// HSA is used to obtain the per-agent properties that HIP does not expose +/// directly (SIMDs per CU, max waves per CU, XCC count). An HSA failure is +/// non-fatal: we fall back to the static `AmdArchDb` presets for those +/// fields. HSA is loaded into HIP's link-map namespace so both share a +/// single KFD session (HIP initialises HSA internally). +struct HsaRuntime { + rocm_loader::LoadedLibrary lib; + hsa_status_t (*init)(void) = nullptr; + hsa_status_t (*iterateAgents)(hsa_status_t (*)(hsa_agent_t, void *), + void *) = nullptr; + hsa_status_t (*agentGetInfo)(hsa_agent_t, hsa_agent_info_t, void *) = nullptr; +}; + +HsaRuntime loadHsaRuntime(void *hipHandle) { + HsaRuntime rt; + rt.lib = rocm_loader::loadRocmLibrary(rocm_loader::Library::Hsa, hipHandle); + if (!rt.lib.handle) { + LLVM_DEBUG(llvm::dbgs() << "rock-amd-arch-db: libhsa-runtime64 not " + "found; HSA-derived fields will use presets\n"); + return rt; + } + + rt.init = reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, "hsa_init")); + rt.iterateAgents = reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, "hsa_iterate_agents")); + rt.agentGetInfo = + reinterpret_cast( + rocm_loader::resolveRocmSymbol(rt.lib, "hsa_agent_get_info")); + + if (!rt.init || !rt.iterateAgents || !rt.agentGetInfo) { + LLVM_DEBUG(llvm::dbgs() + << "rock-amd-arch-db: HSA loaded but required symbols are " + "missing; HSA-derived fields will fall back to presets\n"); + rt.lib.handle = nullptr; + return rt; + } + + // HIP indirectly calls `hsa_init()`; do it explicitly so we can use HSA + // without HIP (e.g. future tests that want HSA-only queries). The runtime + // is reference-counted internally, so double-initialisation is harmless. + if (rt.init() != HSA_STATUS_SUCCESS) { + LLVM_DEBUG(llvm::dbgs() << "rock-amd-arch-db: hsa_init() failed\n"); + rt.lib.handle = nullptr; + } + return rt; +} + +const HsaRuntime &getHsaRuntime() { + static HsaRuntime rt = loadHsaRuntime(getHipRuntime().lib.handle); + return rt; +} + +/// Agent-iteration callback and per-device state for HSA queries. +struct AgentQuery { + const HsaRuntime *hsa; + uint32_t targetDeviceId; int numCpus; - // Output fields: uint32_t simdsPerCU; uint32_t maxWavesPerCU; uint32_t numXCC; + bool found; }; -AmdArchInfo fetchNativeArchInfo(const hipDeviceProp_t &prop, - AgentInfo &agentInfo) { - auto ret = lookupArchInfo(prop.gcnArchName); // get baseline +// Adapted from `rocminfo.cc`; see +// https://github.com/ROCm/rocm-systems/blob/develop/projects/rocminfo/rocminfo.cc +hsa_status_t acquireAgentInfo(hsa_agent_t agent, void *data) { + auto *q = static_cast(data); + const HsaRuntime &hsa = *q->hsa; - checkAndSetInfo("(HIP) minNumCU", ret.minNumCU, prop.multiProcessorCount); - checkAndSetInfo("(HIP) waveSize", ret.waveSize, prop.warpSize); - checkAndSetInfo("(HIP) totalSharedMemPerCU", ret.totalSharedMemPerCU, - prop.maxSharedMemoryPerMultiProcessor); - checkAndSetInfo("(HIP) maxSharedMemPerWG", ret.maxSharedMemPerWG, - prop.sharedMemPerBlock); + hsa_device_type_t deviceType; + if (hsa_status_t err = + hsa.agentGetInfo(agent, HSA_AGENT_INFO_DEVICE, &deviceType); + err != HSA_STATUS_SUCCESS) + return err; + + if (deviceType != HSA_DEVICE_TYPE_GPU) { + ++q->numCpus; + return HSA_STATUS_SUCCESS; + } -// We cannot get those values under Windows, since HSA is not supported. -#ifndef _WIN32 - checkAndSetInfo("(HSA) numEUPerCU", ret.numEUPerCU, agentInfo.simdsPerCU); - checkAndSetInfo("(HSA) maxWavesPerEU", ret.maxWavesPerEU, - agentInfo.maxWavesPerCU / agentInfo.simdsPerCU); - checkAndSetInfo("(HSA) maxNumXCC", ret.maxNumXCC, agentInfo.numXCC); -#endif + uint32_t internalNodeId = 0; + if (hsa_status_t err = hsa.agentGetInfo( + agent, + static_cast(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID), + &internalNodeId); + err != HSA_STATUS_SUCCESS) + return err; + + if (internalNodeId < static_cast(q->numCpus)) + return HSA_STATUS_SUCCESS; + + uint32_t gpuId = internalNodeId - static_cast(q->numCpus); + if (gpuId != q->targetDeviceId) + return HSA_STATUS_SUCCESS; + + if (hsa_status_t err = hsa.agentGetInfo( + agent, + static_cast(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU), + &q->simdsPerCU); + err != HSA_STATUS_SUCCESS) + return err; + if (hsa_status_t err = hsa.agentGetInfo( + agent, + static_cast(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU), + &q->maxWavesPerCU); + err != HSA_STATUS_SUCCESS) + return err; + if (hsa_status_t err = hsa.agentGetInfo( + agent, static_cast(HSA_AMD_AGENT_INFO_NUM_XCC), + &q->numXCC); + err != HSA_STATUS_SUCCESS) + return err; + + q->found = true; + return HSA_STATUS_SUCCESS; +} +#endif // _WIN32 - // TODO: Add missing fields: - // - totalSGPRPerEU - // - totalVGPRPerEU - // - defaultFeatures - // - hasOcpFp8ConversionInstrs - return ret; +//===----------------------------------------------------------------------===// +// Native arch inference. +//===----------------------------------------------------------------------===// + +/// Apply the WGP-as-CU correction for Navi-class GPUs. HIP/HSA report +/// wavefront-32 GPUs in WGP mode, which halves several per-CU metrics +/// compared to our rocMLIR convention; scale them back up. The +/// `sharedMemPerCU` out-parameter is only updated when the caller has +/// matching HSA data; HIP-only paths leave it alone. +void applyNaviCorrection(uint32_t warpSize, uint32_t &simdsPerCU, + uint32_t &maxWavesPerCU, int64_t &sharedMemPerCU) { + if (warpSize != 32) + return; + simdsPerCU *= 2; + maxWavesPerCU *= 2; + sharedMemPerCU *= 2; } -#define RET_IF_HSA_ERR(err) \ - { \ - if ((err) != HSA_STATUS_SUCCESS) { \ - return err; \ - } \ +template +std::enable_if_t, void> +checkAndSetInfo(StringRef name, LHS &lhs, RHS &&rhs) { + if (lhs != static_cast(rhs)) { + LLVM_DEBUG(llvm::dbgs() << "NOTE: Value discrepancy for " << name << ": " + << lhs << " (old) != " << rhs + << " (new). Proceeding with " << rhs << ".\n"); + lhs = std::forward(rhs); } +} -// hsa_iterate_agents expects a callback function (acquireAgentInfo in this -// case) with one void* argument which contains arbitrary data to be used by the -// called function. Each time the callback is invoked, it is called with a -// different HSA agent and the pointer (i.e., the void* argument is shared -// across all calls). That is also why we count the number of CPUs, since we -// need to match the HIP deviceId with the HSA agent index. -// -// See hsa_iterate_agents documentation in -// https://rocm.docs.amd.com/projects/ROCR-Runtime/en/latest/api-reference/api.html -// for more information. -static hsa_status_t acquireAgentInfo(hsa_agent_t agent, void *data) { - // Use HSA to get data not exposed by HIP. - // Based on: - // https://github.com/ROCm/rocm-systems/blob/develop/projects/rocminfo/rocminfo.cc - hsa_status_t err; - AgentInfo *agentI = reinterpret_cast(data); +/// Query HIP (+ HSA where available) for the running device and adjust the +/// static preset that matches `gcnArchName`. Returns `std::nullopt` if HIP +/// cannot be loaded or the query fails. +std::optional tryQueryNativeArchInfo(unsigned deviceId, + std::string &gcnArchName) { + const HipRuntime &hip = getHipRuntime(); + if (!hip.lib.handle) + return std::nullopt; + + hipDeviceProp_t prop{}; + if (hip.getDeviceProperties(&prop, static_cast(deviceId)) != hipSuccess) + return std::nullopt; + gcnArchName = prop.gcnArchName; + LLVM_DEBUG(llvm::dbgs() << "gcnArchName: " << gcnArchName << "\n"); + + // Query HSA up front so we can apply the Navi WGP-as-CU correction to the + // shared-memory figure before it reaches `checkAndSetInfo`. This matches + // the shim's semantics exactly (@31416e7): on Navi-with-HSA the per-CU + // shared memory reported by HIP is doubled to account for WGP mode. If + // HSA is unavailable we skip the correction rather than silently emitting + // a value that differs from the static preset. + uint32_t simdsPerCU = 0; + uint32_t maxWavesPerCU = 0; + uint32_t numXCC = 0; + bool hsaValid = false; + int64_t sharedMemPerCU = + static_cast(prop.maxSharedMemoryPerMultiProcessor); - hsa_device_type_t deviceType; - err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &deviceType); - RET_IF_HSA_ERR(err); - - if (HSA_DEVICE_TYPE_GPU == deviceType) { - // This a GPU, check if its the GPU that we are looking for. - uint32_t internalNodeId; - err = hsa_agent_get_info( - agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_DRIVER_NODE_ID, - &internalNodeId); - RET_IF_HSA_ERR(err); - - unsigned gpuDeviceId = internalNodeId - agentI->numCpus; - - if (gpuDeviceId == agentI->deviceId) { - // This is the GPU that we want to check. - err = hsa_agent_get_info( - agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU, - &agentI->simdsPerCU); - RET_IF_HSA_ERR(err); - - err = hsa_agent_get_info( - agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU, - &agentI->maxWavesPerCU); - RET_IF_HSA_ERR(err); - - err = hsa_agent_get_info( - agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_NUM_XCC, &agentI->numXCC); - RET_IF_HSA_ERR(err); +#ifndef _WIN32 + const HsaRuntime &hsa = getHsaRuntime(); + if (hsa.lib.handle) { + AgentQuery q{}; + q.hsa = &hsa; + q.targetDeviceId = deviceId; + if (hsa.iterateAgents(acquireAgentInfo, &q) == HSA_STATUS_SUCCESS && + q.found && q.simdsPerCU != 0) { + simdsPerCU = q.simdsPerCU; + maxWavesPerCU = q.maxWavesPerCU; + numXCC = q.numXCC; + hsaValid = true; + applyNaviCorrection(static_cast(prop.warpSize), simdsPerCU, + maxWavesPerCU, sharedMemPerCU); + } else { + LLVM_DEBUG(llvm::dbgs() + << "rock-amd-arch-db: HSA agent query for device " << deviceId + << " failed; keeping preset values\n"); } - } else { - agentI->numCpus++; } +#endif - return HSA_STATUS_SUCCESS; -} - -void fixNaviProperties(AgentInfo *agentI, hipDeviceProp_t *prop) { - // Fix per CU metrics in Navi GPUs due to WGPs. - // I wonder why we have to implement this logic instead of relying - // on HIP to do this. - // - // Navi AMD docs define a CU as "One half of a WGP. Contains 2 SIMD32’s that - // share one path to memory" In this context we treat a WGP as CU, so we need - // to double simdsPerCU, totalSharedMemPerCU and - // maxSharedMemoryPerMultiProcessor. This is consistent with the behavior of - // amdgpu target in LLVM. They say: "Per CU" really means "per whatever - // functional block the waves of a workgroup must share" This is also - // mentioned on HIP multiProcessorCount field: "When the GPU works in Compute - // Unit (CU) mode, this value equals the number of CUs; when in Workgroup - // Processor (WGP) mode, this value equels half of CUs, because a single WGP - // contains two CUs" - // - // References: - // - - // https://rocm.docs.amd.com/projects/HIP/en/docs-6.0.2/user_guide/hip_rtc.html#cu-mode-vs-wgp-mode - // - - // https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna3-shader-instruction-set-architecture-feb-2023_0.pdf - // - - // https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp - - // TODO: Can we check WGP mode in a better way instead of checking warp size? - if (prop->warpSize == 32) { - agentI->simdsPerCU *= 2; - agentI->maxWavesPerCU *= 2; - prop->maxSharedMemoryPerMultiProcessor *= 2; + AmdArchInfo ret = lookupArchInfo(gcnArchName); + checkAndSetInfo("(HIP) minNumCU", ret.minNumCU, + static_cast(prop.multiProcessorCount)); + checkAndSetInfo("(HIP) waveSize", ret.waveSize, + static_cast(prop.warpSize)); + checkAndSetInfo("(HIP) totalSharedMemPerCU", ret.totalSharedMemPerCU, + sharedMemPerCU); + checkAndSetInfo("(HIP) maxSharedMemPerWG", ret.maxSharedMemPerWG, + static_cast(prop.sharedMemPerBlock)); + + if (hsaValid) { + checkAndSetInfo("(HSA) numEUPerCU", ret.numEUPerCU, + static_cast(simdsPerCU)); + checkAndSetInfo("(HSA) maxWavesPerEU", ret.maxWavesPerEU, + static_cast(maxWavesPerCU / simdsPerCU)); + checkAndSetInfo("(HSA) maxNumXCC", ret.maxNumXCC, + static_cast(numXCC)); } + + // NOTE: the following AmdArchInfo fields are not yet sourced from hardware + // and therefore keep their static-preset values from `lookupArchInfo` above: + // - totalSGPRPerEU + // - totalVGPRPerEU + // - defaultFeatures + // - hasOcpFp8ConversionInstrs + // Adding HIP/HSA queries for these is tracked as part of the original + // native-arch work (PR #1790). + return ret; } -AmdArchInfo nativeArchInfo(unsigned deviceId = 0) { +AmdArchInfo nativeArchInfo(unsigned deviceId) { + // Cache is keyed by deviceId, NOT by `gcnArchName`. Two GPUs with the same + // `gcnArchName` (e.g. two gfx942 cards) can still report different per-device + // properties (CU count, XCC count, per-CU shared memory on binned variants), + // so caching by arch alone would silently return the first device's data for + // every later device on a same-arch multi-GPU system. The deviceId-keyed + // cache pins each device to the values queried from the actual hardware. static std::mutex m; - static std::unordered_map cache; + static llvm::DenseMap cache; LLVM_DEBUG(llvm::dbgs() << "Retrieving native arch info for device " << deviceId << "...\n"); - hipDeviceProp_t prop; - if (auto err = hipGetDeviceProperties(&prop, deviceId); err != hipSuccess) { - auto reason = "hipGetDeviceProperties failed with error: " + - std::string(hipGetErrorString(err)); - llvm::report_fatal_error(reason.c_str()); + { + std::lock_guard lock(m); + if (auto it = cache.find(deviceId); it != cache.end()) + return it->second; } - LLVM_DEBUG(llvm::dbgs() << "gcnArchName: " << prop.gcnArchName << "\n"); - - AgentInfo agentInfo; -#ifndef _WIN32 - agentInfo.numCpus = 0; - agentInfo.deviceId = deviceId; - hsa_status_t err = hsa_iterate_agents(acquireAgentInfo, &agentInfo); - if (err != HSA_STATUS_SUCCESS) { - char errVal[12]; - const char *errStr = nullptr; - if (hsa_status_string(err, (const char **)&errStr) != HSA_STATUS_SUCCESS) { - snprintf(&(errVal[0]), sizeof(errVal), "%#x", (uint32_t)err); - errStr = &(errVal[0]); - } - llvm::report_fatal_error(errStr); - } - - fixNaviProperties(&agentInfo, &prop); -#endif + std::string gcnArchName; + std::optional queried = + tryQueryNativeArchInfo(deviceId, gcnArchName); + if (!queried) + llvm::report_fatal_error( + "Failed to query AMD GPU arch runtime for native architecture " + "detection. Ensure a ROCm installation with libamdhip64 is visible " + "via LD_LIBRARY_PATH / RPATH, and that the requested device id is " + "valid."); std::lock_guard lock(m); - - auto it = cache.find(prop.gcnArchName); - if (it == cache.end()) { - LLVM_DEBUG(llvm::dbgs() << "Cache miss! Fetching native arch info...\n"); - it = cache.emplace(prop.gcnArchName, fetchNativeArchInfo(prop, agentInfo)) - .first; + auto [it, inserted] = cache.try_emplace(deviceId, *queried); + if (inserted) { + LLVM_DEBUG(llvm::dbgs() << "Cache miss for device " << deviceId + << " (gcnArchName=" << gcnArchName << ")\n"); } - return it->second; } } // anonymous namespace -#endif // _WIN32 - AmdArchInfo mlir::rock::lookupArchInfo(StringRef arch) { // Keep this implementation in sync with // mlir/test/lit.site.cfg.py.in:set_arch_features() auto [chip, deviceId] = parseArchString(arch); - if (chip == "native") { -#ifdef _WIN32 - llvm_unreachable("native arch lookup is not supported on Windows"); -#else + if (chip == "native") return nativeArchInfo(deviceId); -#endif - } StringRef minor = chip.take_back(2); StringRef major = chip.slice(0, chip.size() - 2); if (major == "gfx9") { @@ -396,6 +568,28 @@ AmdArchInfo mlir::rock::lookupArchInfo(StringRef arch) { llvm_unreachable(msg.c_str()); } +unsigned mlir::rock::nativeDeviceCount() { + const HipRuntime &hip = getHipRuntime(); + if (!hip.lib.handle) + return 0; + int count = 0; + if (hip.getDeviceCount(&count) != hipSuccess) + return 0; + if (count < 0) + return 0; + return static_cast(count); +} + +std::string mlir::rock::nativeArchName(unsigned deviceId) { + const HipRuntime &hip = getHipRuntime(); + if (!hip.lib.handle) + return std::string(); + hipDeviceProp_t prop{}; + if (hip.getDeviceProperties(&prop, static_cast(deviceId)) != hipSuccess) + return std::string(); + return std::string(prop.gcnArchName); +} + GemmFeatures mlir::rock::AmdArchInfo::getDefaultFeatures(Type dataType) { GemmFeatures theseFeatures = defaultFeatures; bool isWmma = bitEnumContainsAll(theseFeatures, GemmFeatures::wmma); diff --git a/mlir/lib/Dialect/Rock/IR/CMakeLists.txt b/mlir/lib/Dialect/Rock/IR/CMakeLists.txt index b15e2aea59f5..5e69ca31398c 100644 --- a/mlir/lib/Dialect/Rock/IR/CMakeLists.txt +++ b/mlir/lib/Dialect/Rock/IR/CMakeLists.txt @@ -45,24 +45,58 @@ target_link_libraries(MLIRRockOps MLIRROCDLDialect MLIRSupport MLIRLinalgTransformOps + PRIVATE + MLIRRocmRuntimeLoader ) -# ROCm on Windows does not support HSA. -# Also HIP is not supported on our CI, -# so we cannot use either of them. -if (NOT WIN32) - # From: https://github.com/ROCm/rocMLIR/blob/b6c726a324d2807494f770ed915af0e54cc6cb3f/external/llvm-project/mlir/lib/ExecutionEngine/CMakeLists.txt#L414 - find_package(hip REQUIRED PATHS ${ROCM_PATH}) - # We do need both of these. One is for linking, to get the HIP - # functions. The other is for compiling, to get -D__HIP_PLATFORM_AMD__ - target_link_libraries(MLIRRockOps PUBLIC - hip::host hip::amdhip64 - ) - target_link_libraries(obj.MLIRRockOps PUBLIC - hip::host hip::amdhip64 - ) +# MLIRRockOps deliberately has no link-time dependency on `libamdhip64` / +# `libhsa-runtime64`. `AmdArchDb.cpp` needs the HIP and HSA headers to know +# the POD layout of `hipDeviceProp_t` and a handful of HSA enums, but it +# resolves the actual functions at runtime via `dlopen`/`LoadLibraryW`. +# Avoiding the link-time dependency keeps ROCm's `libamd_comgr` and its +# embedded `libLLVM.so.` out of every binary that links `MLIRRockOps`; +# otherwise the two LLVM copies race for the `cl::opt` global registry at +# static-init time and abort the process. +# +# The loader (`dlopen` / `LoadLibraryW`) lets the dynamic linker find the +# runtime libraries via the standard search path (`LD_LIBRARY_PATH` / RPATH +# on POSIX, the DLL search order on Windows); no compile-time install path +# is baked in. +find_package(hip QUIET) +if (NOT hip_FOUND) + message(FATAL_ERROR + "MLIRRockOps needs the HIP headers to describe `hipDeviceProp_t` for " + "the `rock.arch = \"native\"` delay-loader in AmdArchDb.cpp. Install " + "the ROCm `hip` CMake package (e.g. set `CMAKE_PREFIX_PATH=/opt/rocm`).") +endif() +# Apply to both the object library (the actual compile target) and the +# aggregate shared library, for parity with how upstream ExecutionEngine +# wires include dirs for its HIP-touching TUs. +foreach(_t MLIRRockOps obj.MLIRRockOps) + if (TARGET ${_t}) + target_include_directories(${_t} SYSTEM PRIVATE ${hip_INCLUDE_DIRS}) + endif() +endforeach() +unset(_t) - # From: https://github.com/ROCm/rocm-systems/blob/c53bdb9643afc2ef943fd25113a1c704d39aa6b4/projects/rocminfo/CMakeLists.txt - find_package(hsa-runtime64 1.0 REQUIRED) - target_link_libraries(MLIRRockOps PUBLIC hsa-runtime64::hsa-runtime64) +if (NOT WIN32) + find_package(hsa-runtime64 QUIET) + if (hsa-runtime64_FOUND) + get_target_property(_hsa_include_dirs hsa-runtime64::hsa-runtime64 + INTERFACE_INCLUDE_DIRECTORIES) + if (_hsa_include_dirs) + foreach(_t MLIRRockOps obj.MLIRRockOps) + if (TARGET ${_t}) + target_include_directories(${_t} SYSTEM PRIVATE + ${_hsa_include_dirs}) + endif() + endforeach() + unset(_t) + endif() + unset(_hsa_include_dirs) + else() + message(STATUS + "hsa-runtime64 CMake package not found; AmdArchDb.cpp will look for " + "HSA headers on the default include path.") + endif() endif() diff --git a/mlir/lib/Dialect/Rock/Tuning/ParamLookupTable.cpp b/mlir/lib/Dialect/Rock/Tuning/ParamLookupTable.cpp index 2dad1515231f..7a3733ae804f 100644 --- a/mlir/lib/Dialect/Rock/Tuning/ParamLookupTable.cpp +++ b/mlir/lib/Dialect/Rock/Tuning/ParamLookupTable.cpp @@ -3,8 +3,12 @@ #include "mlir/Dialect/Rock/Tuning/GridwiseGemmGemmParams.h" #include "mlir/Dialect/Rock/Tuning/GridwiseGemmParams.h" #include "llvm/ADT/StringExtras.h" +#include "llvm/ADT/StringMap.h" #include "llvm/Support/Debug.h" +#include +#include + #define DEBUG_TYPE "rock-tuning-parameter" using namespace mlir; @@ -87,11 +91,47 @@ ParamLookupTable::getRelatives(StringRef target) { template StringRef ParamLookupTable::normalizeArch(StringRef arch) { + // Resolve `native[:N]` to the hardware-reported arch name via the same + // runtime that `lookupArchInfo` uses. The resolved string is cached in a + // process-wide map so the returned `StringRef` stays valid for the lifetime + // of the process (required by the `static const std::map` + // lookup table this value feeds into). + if (arch.consume_front("native")) { + static std::mutex m; + static llvm::StringMap resolvedCache; + std::lock_guard lock(m); + // Re-form the original key for the cache lookup. + std::string cacheKey = ("native" + arch).str(); + auto [it, inserted] = resolvedCache.try_emplace(cacheKey); + if (inserted) { + unsigned deviceId = 0; + if (arch.consume_front(":")) { + // Parse must succeed and fit in `unsigned`. Silently treating + // `native:foo` or `native:` as `native:0` would mask user error + // and target the wrong GPU on multi-GPU systems. + unsigned long long parsed = 0; + if (arch.empty() || llvm::getAsUnsignedInteger(arch, 0, parsed) || + parsed > std::numeric_limits::max()) + llvm::report_fatal_error( + Twine("Invalid `") + cacheKey + + "`: the suffix after `native:` must be a non-negative integer " + "device id (got `" + + arch + "`)."); + deviceId = static_cast(parsed); + } + it->second = nativeArchName(deviceId); + if (it->second.empty()) + llvm::report_fatal_error( + Twine("Failed to resolve `") + cacheKey + + "`: AMD GPU arch runtime unavailable or device not present"); + } + return normalizeArch(it->second); + } + auto gfxPos = arch.find("gfx"); - if (gfxPos == StringRef::npos) { + if (gfxPos == StringRef::npos) llvm::report_fatal_error(Twine("Invalid architecture string: ") + arch); - } - auto remaining = arch.substr(gfxPos); + StringRef remaining = arch.substr(gfxPos); auto endPos = remaining.find_if_not([](char c) { return llvm::isAlnum(c); }, 3); return remaining.substr(0, endPos); diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt index feecf3a7769c..4b950dbb2ce6 100644 --- a/mlir/lib/ExecutionEngine/CMakeLists.txt +++ b/mlir/lib/ExecutionEngine/CMakeLists.txt @@ -4,6 +4,7 @@ if (BUILD_FAT_LIBROCKCOMPILER) endif() add_mlir_library(conv-validation-wrappers SHARED EXCLUDE_FROM_LIBMLIR + PARTIAL_SOURCES_INTENDED ${disable_install} conv-validation-wrappers.cpp ) @@ -20,4 +21,3 @@ if (NOT MSVC) set_source_files_properties(conv-validation-wrappers.cpp PROPERTIES COMPILE_OPTIONS "-Wno-gnu-anonymous-struct;-Wno-nested-anon-types;-Wno-return-type-c-linkage;-Wno-c++98-compat-extra-semi") endif() - diff --git a/mlir/test/Dialect/Rock/Loader/check_dynsym_only_mgpu.sh b/mlir/test/Dialect/Rock/Loader/check_dynsym_only_mgpu.sh new file mode 100755 index 000000000000..719e371239b2 --- /dev/null +++ b/mlir/test/Dialect/Rock/Loader/check_dynsym_only_mgpu.sh @@ -0,0 +1,75 @@ +#!/usr/bin/env bash +# Helper for `dynsym_only_mgpu.test`. Asserts that `libmlir_rocm_runtime.so` +# exports nothing but the `mgpu*` C entry points. Skips when the library is +# not built (`MLIR_ENABLE_ROCM_RUNNER=OFF`) or when `nm` is not available. +# +# Usage: check_dynsym_only_mgpu.sh +# +# Version-agnostic: we glob `libmlir_rocm_runtime.so*` (preferring the SONAME- +# versioned file, falling back to the unversioned dev symlink), so a future +# LLVM bump from `.so.23.0git` to `.so.24.0git` does not silently turn the +# test into a no-op. + +set -u +shlib_dir="${1:?shlib dir}" + +# Pick the SONAME-versioned file first (what runtime consumers actually +# `dlopen`); fall back to the dev symlink. The version suffix glob is +# anchored on a digit so we never accidentally match `.so..dwo` +# (split-DWARF), `.debug`, or `.dbg` companion files. `nullglob` makes +# the glob expand to nothing on a clean miss instead of leaving the +# literal pattern. (Brackets must stay outside double-quotes for bash +# to treat them as a character class.) +shopt -s nullglob +candidates=("${shlib_dir}"/libmlir_rocm_runtime.so.[0-9]* \ + "${shlib_dir}/libmlir_rocm_runtime.so") +shopt -u nullglob + +target="" +for cand in "${candidates[@]}"; do + case "${cand}" in + *.dwo|*.debug|*.dbg) continue ;; + esac + if [ -f "${cand}" ] || [ -L "${cand}" ]; then + target="${cand}" + break + fi +done + +if [ -z "${target}" ]; then + echo "dynsym_only_mgpu: skipping; libmlir_rocm_runtime.so* not built." >&2 + exit 0 +fi +if ! command -v nm >/dev/null 2>&1; then + echo "dynsym_only_mgpu: skipping; nm is not available." >&2 + exit 0 +fi + +# Capture `nm` output separately from the awk filter so a `nm` failure +# surfaces immediately rather than being silently swallowed. +if ! nm_out="$(nm -D --defined-only "${target}" 2>/dev/null)"; then + echo "FAIL: nm -D --defined-only ${target} failed" >&2 + exit 1 +fi + +# Single `awk` pass: drop linker pseudo-symbols, partition into "mgpu*" +# (allowed) and everything else (forbidden). Tag the two categories with +# `OK ` / `BAD ` line prefixes so the caller can split them apart with a +# single `grep` per category without sentinel lines or empty spacers. +report="$(awk ' + $3 ~ /^(_init|_fini|_edata|_end|__bss_start)$/ { next } + $3 == "" { next } + $3 ~ /^mgpu/ { print "OK " $3; next } + { print "BAD " $3 }' <<<"${nm_out}")" + +bad="$(grep '^BAD ' <<<"${report}" | cut -d' ' -f2- || true)" +if [ -n "${bad}" ]; then + echo "FAIL: ${target} exports forbidden non-mgpu symbols:" >&2 + while IFS= read -r sym; do + echo " ${sym}" >&2 + done <<<"${bad}" + exit 1 +fi + +ok="$(grep -c '^OK ' <<<"${report}" || true)" +echo "dynsym_only_mgpu: ${target}: clean (${ok} mgpu* symbols)." diff --git a/mlir/test/Dialect/Rock/Loader/check_no_rocm_neededs.sh b/mlir/test/Dialect/Rock/Loader/check_no_rocm_neededs.sh new file mode 100755 index 000000000000..b175d4e5f1bf --- /dev/null +++ b/mlir/test/Dialect/Rock/Loader/check_no_rocm_neededs.sh @@ -0,0 +1,101 @@ +#!/usr/bin/env bash +# Helper for `no_rocm_neededs.test`. Scans build artefacts and asserts that +# none of them transitively `NEED` a ROCm runtime library. +# +# Usage: check_no_rocm_neededs.sh +# +# Version-agnostic by design: the artefact list uses globs against library +# *base names*, never hardcoded `.so..` suffixes, so a future +# LLVM / rocMLIR version bump does not silently turn this test into a no-op. + +set -u +shlib_dir="${1:?shlib dir}" +tools_dir="${2:?tools dir}" + +# SONAMEs that must never appear in `NEEDED`. We anchor on the full library +# basename (`.so` / `.dll`) followed either by end-of-string or by +# a version separator (`.` / `-git`). This rejects false +# positives like `libamd_comgr_helper.so.1` (extra suffix on the basename) or +# the in-tree `libLLVMSupport.so.23.0git` (component-decorated `libLLVM*`), +# while still matching every real ROCm SONAME we have ever seen +# (`libamdhip64.so`, `libamdhip64.so.7`, `libhiprtc.so.7`, +# `libamd_comgr.so.3`, `libLLVM.so.23.0git`, `libLLVM-23git.so` -- the latter +# normalised to its `libLLVM.so.` SONAME by the dynamic linker). +forbidden_re='^lib(amdhip64|hiprtc|amd_comgr|LLVM)\.(so|dll)([.-]|$)' + +# Resolve the artefact set at run time. Tools have no extension and are added +# only when they actually exist; shared libraries are matched by base name + +# wildcard suffix so any version-decorated variant is picked up. A run with +# zero artefacts is treated as a wrong invocation and fails loudly so the test +# never silently degrades to a no-op (e.g. when the caller passes the wrong +# build directory). +shopt -s nullglob +artefacts=() +for t in rocmlir-driver rocmlir-opt rocmlir-gen rocmlir-tuning-driver \ + xmir-runner mlir-runner; do + if [ -e "${tools_dir}/${t}" ]; then + artefacts+=("${tools_dir}/${t}") + fi +done +# Match `.so....` and the bare `.so` symlink, but NOT side files +# like `.so.<...>.dwo` (split-DWARF debug info), `.so.<...>.debug` +# (separate debug info), or `.so.<...>.dbg`. We anchor the version suffix +# on a digit so non-version decorations are rejected; the trailing case +# statement also drops any debug companion that slipped through. +# (Brackets must stay outside double-quotes for bash to treat them as a +# character class.) +for g in libMLIRRockOps libMLIRRocmRuntimeLoader libmlir_rocm_runtime \ + libMLIRRocmExecutionEngineUtils; do + for f in "${shlib_dir}/${g}".so.[0-9]* "${shlib_dir}/${g}".so; do + case "${f}" in + *.dwo|*.debug|*.dbg) continue ;; + esac + if [ -f "${f}" ] || [ -L "${f}" ]; then + artefacts+=("${f}") + fi + done +done +shopt -u nullglob + +if [ "${#artefacts[@]}" -eq 0 ]; then + echo "no_rocm_neededs: no artefacts found under" \ + "tools=${tools_dir} shlib=${shlib_dir}" >&2 + echo "(this usually means the caller passed the wrong build directory;" \ + "check the lit substitutions \`%rocmlir_tools_dir\` and" \ + "\`%rocmlir_shlib_dir\`.)" >&2 + exit 1 +fi +if ! command -v readelf >/dev/null 2>&1; then + echo "no_rocm_neededs: skipping; readelf is not available" >&2 + exit 0 +fi + +# Per artefact: get `readelf -d`, split out the NEEDED entries, then grep for +# any forbidden basename. We capture `readelf` and `awk` outputs separately +# from the final `grep` so a failure in either tool surfaces as a non-zero +# exit code and a diagnostic, rather than being silently swallowed by an +# empty pipeline. +failed=0 +for art in "${artefacts[@]}"; do + if ! readelf_out="$(readelf -d "${art}" 2>/dev/null)"; then + echo "FAIL: readelf -d ${art} failed" >&2 + failed=1 + continue + fi + needed="$(awk '/\(NEEDED\)/{ gsub(/[][]/,"",$5); print $5 }' \ + <<<"${readelf_out}")" + bad="$(grep -E "${forbidden_re}" <<<"${needed}" || true)" + if [ -n "${bad}" ]; then + while IFS= read -r soname; do + echo "FAIL: ${art} declares NEEDED ${soname}" >&2 + done <<<"${bad}" + failed=1 + fi +done +checked="${#artefacts[@]}" + +if [ "${failed}" -ne 0 ]; then + echo "no_rocm_neededs: at least one forbidden NEEDED entry was found." >&2 + exit 1 +fi +echo "no_rocm_neededs: ${checked} artefact(s) checked, all clean." diff --git a/mlir/test/Dialect/Rock/Loader/dynsym_only_mgpu.test b/mlir/test/Dialect/Rock/Loader/dynsym_only_mgpu.test new file mode 100644 index 000000000000..efd2623429e0 --- /dev/null +++ b/mlir/test/Dialect/Rock/Loader/dynsym_only_mgpu.test @@ -0,0 +1,18 @@ +# Confirm that `libmlir_rocm_runtime.so` exports only the `mgpu*` C ABI +# entry points required by upstream MLIR's GPU lowering. Anything else +# (notably any LLVM internal such as `llvm::EnableABIBreakingChecks`) +# leaking into the dynsym would re-introduce ODR collisions when the +# library is dlopen-ed alongside an LLVM-using host process. +# +# This is enforced at link time by the `mlir_rocm_runtime.map` version +# script. The test exists to catch accidental regressions when someone +# adds a new source file or a new transitive LLVM dependency. +# +# Version-agnostic: the helper script globs `libmlir_rocm_runtime.so*` +# rather than referencing a specific `.so..git` filename, +# so a future LLVM merge does not silently turn this test into a +# (skipped) no-op. + +# REQUIRES: linux + +# RUN: %S/check_dynsym_only_mgpu.sh %rocmlir_shlib_dir diff --git a/mlir/test/Dialect/Rock/Loader/no_rocm_neededs.test b/mlir/test/Dialect/Rock/Loader/no_rocm_neededs.test new file mode 100644 index 000000000000..ca8bfc7be787 --- /dev/null +++ b/mlir/test/Dialect/Rock/Loader/no_rocm_neededs.test @@ -0,0 +1,27 @@ +# Confirm that the rocMLIR-side and upstream-MLIR-side ROCm-aware artefacts +# never have ROCm runtime libraries (libamdhip64, libhiprtc, libamd_comgr) +# nor ROCm's monolithic libLLVM.so on their dynamic-link `NEEDED` set. +# +# Direct linkage to any of those would defeat the entire delay-load design: +# the runtime loader pulls them in at startup via DT_NEEDED, runs their +# static constructors, and re-introduces the multi-LLVM `cl::opt` collision +# we worked around with `dlmopen(LM_ID_NEWLM, ...)`. +# +# We deliberately scan binaries we ship (rocmlir-driver / -opt / -gen, +# xmir-runner, rocmlir-tuning-driver, mlir-runner) and the ROCm-aware +# shared libraries that compose them. `libMLIRRockOps.so` may legally +# carry `libMLIRRocmRuntimeLoader.so*` (our delay-load helper, which +# itself only depends on LLVMSupport); it must NOT carry `libamdhip64`. +# +# Notes: +# - The helper script parses `readelf -d` output via awk + grep. +# - We accept missing artefacts (e.g. mlir_rocm_runtime is only built +# when `MLIR_ENABLE_ROCM_RUNNER=ON`); we only check what exists. +# - The check matches base names plus a real version separator, so a +# future ROCm major version (e.g. `libamdhip64.so.8`) is still +# caught and a false-positive prefix (e.g. `libamd_comgr_helper.so`) +# is not. + +# REQUIRES: linux + +# RUN: %S/check_no_rocm_neededs.sh %rocmlir_shlib_dir %rocmlir_tools_dir diff --git a/mlir/test/Dialect/Rock/native_arch.mlir b/mlir/test/Dialect/Rock/native_arch.mlir new file mode 100644 index 000000000000..26ed0b6c87e3 --- /dev/null +++ b/mlir/test/Dialect/Rock/native_arch.mlir @@ -0,0 +1,59 @@ +// Verify that `rock.arch = "native"` and `rock.arch = "native:N"` both flow +// through the rock pipeline by resolving to the hardware-reported gfxXXX via +// a delay-loaded HIP runtime, producing the same lowered IR as a kernel +// pinned to the concrete arch. +// +// `--arch native[:N]` requires: +// 1. an AMD GPU visible to the HIP runtime (skipped otherwise via REQUIRES); +// 2. `libamdhip64` present on the dynamic-loader search path -- any +// standard ROCm install satisfies this. The HIP runtime is opened in +// its own link-map namespace by `mlir::rocm_loader::loadRocmLibrary`, +// so this test does NOT pull ROCm's libLLVM into the rocmlir-opt +// process and is therefore safe to run in a build that ships an +// embedded LLVM. + +// REQUIRES: amd-gpu-present + +// RUN: rocmlir-opt -mlir-print-local-scope -rock-affix-params %s | FileCheck %s + +// CHECK-LABEL: @rock_conv_native +// CHECK-SAME: rock.arch = "native" +// CHECK: rock.conv +// CHECK-SAME: params = #rock.general_gemm_params +func.func @rock_conv_native(%filter : memref<1x128x8x3x3xf32>, + %input : memref<128x1x8x32x32xf32>, + %output : memref<128x1x128x30x30xf32>) + attributes {rock.arch = "native"} { + rock.conv(%filter, %input, %output) features = none { + filter_layout = ["g", "k", "c", "0", "1"], + input_layout = ["ni", "gi", "ci", "0i", "1i"], + output_layout = ["no", "go", "ko", "0o", "1o"], + dilations = [1 : index, 1 : index], + strides = [1 : index, 1 : index], + padding = [0 : index, 0 : index, 0 : index, 0 : index] + } : memref<1x128x8x3x3xf32>, memref<128x1x8x32x32xf32>, memref<128x1x128x30x30xf32> + return +} + +// `native:0` selects device #0 explicitly. The lowered params must match the +// `native` (no-suffix) form because `nativeArchName(0)` is the same query +// path that `nativeArchName()` uses with the default device. + +// CHECK-LABEL: @rock_conv_native_device0 +// CHECK-SAME: rock.arch = "native:0" +// CHECK: rock.conv +// CHECK-SAME: params = #rock.general_gemm_params +func.func @rock_conv_native_device0(%filter : memref<1x128x8x3x3xf32>, + %input : memref<128x1x8x32x32xf32>, + %output : memref<128x1x128x30x30xf32>) + attributes {rock.arch = "native:0"} { + rock.conv(%filter, %input, %output) features = none { + filter_layout = ["g", "k", "c", "0", "1"], + input_layout = ["ni", "gi", "ci", "0i", "1i"], + output_layout = ["no", "go", "ko", "0o", "1o"], + dilations = [1 : index, 1 : index], + strides = [1 : index, 1 : index], + padding = [0 : index, 0 : index, 0 : index, 0 : index] + } : memref<1x128x8x3x3xf32>, memref<128x1x8x32x32xf32>, memref<128x1x128x30x30xf32> + return +} diff --git a/mlir/test/Dialect/Rock/native_arch_invalid.mlir b/mlir/test/Dialect/Rock/native_arch_invalid.mlir new file mode 100644 index 000000000000..e46021048bd7 --- /dev/null +++ b/mlir/test/Dialect/Rock/native_arch_invalid.mlir @@ -0,0 +1,34 @@ +// Confirm that malformed `rock.arch = "native:..."` strings produce a fatal +// error, rather than silently falling back to device 0. The latter behaviour +// (which we explicitly fixed) would have masked user typos and silently +// targeted the wrong GPU on multi-GPU systems. +// +// `parseArchString` reports the failure via `llvm::report_fatal_error`, +// which `abort()`s the process; we use `not --crash` to match that signal- +// terminated exit code (plain `not` only matches non-signal non-zero exits). +// The fatal error fires before any HIP query, so this test does NOT need a +// real GPU and can run anywhere. + +// REQUIRES: linux + +// RUN: not --crash rocmlir-opt -mlir-print-local-scope -rock-affix-params %s 2>&1 \ +// RUN: | FileCheck %s + +// CHECK: LLVM ERROR +// CHECK-SAME: native:foo +// CHECK-SAME: must be a non-negative integer device id + +func.func @rock_conv_native_invalid(%filter : memref<1x128x8x3x3xf32>, + %input : memref<128x1x8x32x32xf32>, + %output : memref<128x1x128x30x30xf32>) + attributes {rock.arch = "native:foo"} { + rock.conv(%filter, %input, %output) features = none { + filter_layout = ["g", "k", "c", "0", "1"], + input_layout = ["ni", "gi", "ci", "0i", "1i"], + output_layout = ["no", "go", "ko", "0o", "1o"], + dilations = [1 : index, 1 : index], + strides = [1 : index, 1 : index], + padding = [0 : index, 0 : index, 0 : index, 0 : index] + } : memref<1x128x8x3x3xf32>, memref<128x1x8x32x32xf32>, memref<128x1x128x30x30xf32> + return +} diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py index ecf63326d695..a4bb859c5e98 100644 --- a/mlir/test/lit.cfg.py +++ b/mlir/test/lit.cfg.py @@ -21,7 +21,7 @@ config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell) # suffixes: A list of file extensions to treat as test files. -config.suffixes = ['.td', '.mlir', '.toy', '.ll', '.tc', '.py'] +config.suffixes = ['.td', '.mlir', '.toy', '.ll', '.tc', '.py', '.test'] # test_source_root: The root path where tests are located. config.test_source_root = os.path.dirname(__file__) @@ -32,6 +32,8 @@ config.substitutions.append(('%PATH%', config.environment['PATH'])) config.substitutions.append(('%shlibext', config.llvm_shlib_ext)) config.substitutions.append(("%mlir_src_root", config.mlir_src_root)) +config.substitutions.append(('%rocmlir_shlib_dir', config.rocmlir_shlib_dir)) +config.substitutions.append(('%rocmlir_tools_dir', config.mlir_rock_tools_dir)) config.substitutions.append(('%random_data', config.random_data)) config.substitutions.append( ('%constrained_float_range_random_data', config.constrained_float_range_random_data)) @@ -73,7 +75,11 @@ # directories. config.excludes = [ 'Inputs', 'CMakeLists.txt', 'README.txt', 'LICENSE.txt', 'lit.cfg.py', 'lit.site.cfg.py', - 'common_utils' + 'common_utils', + # `.sh` helpers invoked from `.test` files via `%S/check_*.sh`. They are + # not tests themselves (lit's default suffixes don't include `.sh`, so + # this is more belt-and-braces than required). + 'check_no_rocm_neededs.sh', 'check_dynsym_only_mgpu.sh', ] # test_source_root: The root path where tests are located. diff --git a/mlir/test/lit.site.cfg.py.in b/mlir/test/lit.site.cfg.py.in index fa8e08a52393..cb33dbeb38ce 100644 --- a/mlir/test/lit.site.cfg.py.in +++ b/mlir/test/lit.site.cfg.py.in @@ -113,6 +113,56 @@ if config.rocm_path: except subprocess.CalledProcessError: config.no_AMD_GPU = True +# Tests that exercise the `--arch native:N` runtime path can opt in via +# `// REQUIRES: amd-gpu-present`. The feature is set whenever lit was able to +# enumerate at least one AMD GPU through the HIP runtime. +if not config.no_AMD_GPU: + config.available_features.add('amd-gpu-present') + +# Symbol-contract / NEEDED-cleanliness tests require ELF + readelf and so are +# Linux-only; we expose `linux` as an opt-in feature for them. +if platform.system() == 'Linux': + config.available_features.add('linux') + +# So tests can locate the rocMLIR shared-library output directory (where +# libMLIRRocmRuntimeLoader, libmlir_rocm_runtime, libMLIRRockOps, ... +# live). Resolution order: +# +# 1. Parse the `-Wl,-rpath,` set by `add_mlir_library` out of +# `host_ldflags`. This is the canonical answer in any Ninja build. +# 2. Probe four well-known candidate directories +# (`/lib`, `/lib64`, +# `/external/llvm-project/lib`, `/lib`) and +# pick the first that exists. Covers fat-lib / multi-config +# generators where the rpath may be empty, and `lib` vs `lib64` +# distros. +# +# We never reference a specific LLVM major version, so a future LLVM +# bump or rocMLIR version bump leaves this resolver intact. +import re as _re +def _resolve_rocmlir_shlib_dir(): + # Handle both rpath spellings: + # 1. `-Wl,-rpath -Wl,/path/to/lib` (the form GCC/clang typically + # emit on Linux when CMake's `INSTALL_RPATH` is materialised). + # 2. `-Wl,-rpath,/path/to/lib` (the comma-form, common on macOS + # and from some clang invocations). + for pattern in (r'-Wl,-rpath\s*[-\s]*Wl,([^\s,]+)', + r'-Wl,-rpath,([^\s,]+)'): + m = _re.search(pattern, config.host_ldflags) + if m: + return m.group(1) + for cand in ( + os.path.join(config.mlir_obj_root, "lib"), + os.path.join(config.mlir_obj_root, "lib64"), + os.path.join(config.mlir_obj_root, "external", "llvm-project", "lib"), + os.path.join(config.llvm_obj_root, "lib"), + ): + if os.path.isdir(cand): + return cand + return os.path.join(config.mlir_obj_root, "lib") + +config.rocmlir_shlib_dir = _resolve_rocmlir_shlib_dir() + import lit.llvm lit.llvm.initialize(lit_config, config) diff --git a/mlir/tools/rocmlir-lib/CMakeLists.txt b/mlir/tools/rocmlir-lib/CMakeLists.txt index 2682f74770fa..20c8894e69c9 100644 --- a/mlir/tools/rocmlir-lib/CMakeLists.txt +++ b/mlir/tools/rocmlir-lib/CMakeLists.txt @@ -125,6 +125,12 @@ if(BUILD_FAT_LIBROCKCOMPILER) rocm_install(FILES ${full_output_path} DESTINATION lib) + # The fat archive does not link `libamdhip64` / `libhsa-runtime64`. + # `MLIRRockOps` delay-loads them at run time via dlopen/LoadLibraryW + # (see `mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp`). Downstream consumers + # (MIOpen, MIGraphX) that need `rock.arch = "native"` must have a ROCm + # install visible on the dynamic-loader search path. + # Backward compatibility to provide `librockCompiler` build target. # INTERFACE libraries are not exposed as build targets. add_custom_target(${__library_name} ALL DEPENDS ${LIBRARY_NAME}) diff --git a/mlir/tools/rocmlir-tuning-driver/CMakeLists.txt b/mlir/tools/rocmlir-tuning-driver/CMakeLists.txt index ca392984bae6..d4d55405bf62 100644 --- a/mlir/tools/rocmlir-tuning-driver/CMakeLists.txt +++ b/mlir/tools/rocmlir-tuning-driver/CMakeLists.txt @@ -1,5 +1,7 @@ set(LLVM_OPTIONAL_SOURCES rocmlir-tuning-driver.cpp + CacheFlush.cpp + HipDelayLoad.cpp ) if(MLIR_ENABLE_ROCM_RUNNER OR ROCMLIR_BUILD_TUNING_DRIVER) @@ -17,9 +19,23 @@ set(LIBS add_rocmlir_tool(rocmlir-tuning-driver rocmlir-tuning-driver.cpp CacheFlush.cpp + HipDelayLoad.cpp ) -# Grab HIP again, since we'll be using it directly +# `rocmlir-tuning-driver` does NOT link `libamdhip64` / `libhiprtc` at +# build time. Doing so transitively pulls in `libamd_comgr` and ROCm's +# monolithic `libLLVM.so.`, which fights with rocMLIR's own +# embedded LLVM at process start (`cl::SubCommand` static-init +# collision -> SmallPtrSet "Bucket < End" / "Option already exists"). +# `HipDelayLoad.{h,cpp}` resolves every hipXXX/hiprtcXXX entry point at +# run time via dlmopen, coordinating with RocmSystemDetect through +# `mlirRocmSystemDetectGetHipHandle` so we share a single HSA session +# per process (KFD only allows one). +# +# We still need the HIP and HIPRTC headers (for type definitions like +# `hipDeviceProp_t`, `hipModule_t`, `hiprtcResult`, ...) and the +# `__HIP_PLATFORM_AMD__` define. Pull both from the standard `find_package` +# variables without consuming the imported targets' link components. if (NOT DEFINED ROCM_PATH) if (NOT DEFINED ENV{ROCM_PATH}) set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCm has been installed") @@ -37,39 +53,43 @@ find_package(hip REQUIRED) find_package(hiprtc REQUIRED) set(CMAKE_PREFIX_PATH "${REAL_CMAKE_PREFIX_PATH}") -# Suppress compiler warnings from HIP headers (only when rocm runner is built) -if(MLIR_ENABLE_ROCM_RUNNER) -check_cxx_compiler_flag(-Wno-c++98-compat-extra-semi - CXX_SUPPORTS_NO_CXX98_COMPAT_EXTRA_SEMI_FLAG) -if (CXX_SUPPORTS_CXX98_COMPAT_EXTRA_SEMI_FLAG) - target_compile_options(mlir_rocm_runtime PRIVATE - "-Wno-c++98-compat-extra-semi") -endif() -check_cxx_compiler_flag(-Wno-return-type-c-linkage - CXX_SUPPORTS_WNO_RETURN_TYPE_C_LINKAGE_FLAG) -if (CXX_SUPPORTS_WNO_RETURN_TYPE_C_LINKAGE_FLAG) - target_compile_options(mlir_rocm_runtime PRIVATE - "-Wno-return-type-c-linkage") -endif() -check_cxx_compiler_flag(-Wno-nested-anon-types - CXX_SUPPORTS_WNO_NESTED_ANON_TYPES_FLAG) -if (CXX_SUPPORTS_WNO_NESTED_ANON_TYPES_FLAG) - target_compile_options(mlir_rocm_runtime PRIVATE - "-Wno-nested-anon-types") -endif() -check_cxx_compiler_flag(-Wno-gnu-anonymous-struct - CXX_SUPPORTS_WNO_GNU_ANONYMOUS_STRUCT_FLAG) -if (CXX_SUPPORTS_WNO_GNU_ANONYMOUS_STRUCT_FLAG) - target_compile_options(mlir_rocm_runtime PRIVATE - "-Wno-gnu-anonymous-struct") +target_include_directories(rocmlir-tuning-driver SYSTEM PRIVATE + ${hip_INCLUDE_DIR}) +target_compile_definitions(rocmlir-tuning-driver PRIVATE + __HIP_PLATFORM_AMD__=1) +# `find_package(hiprtc)` exposes its include path either via the imported +# target's INTERFACE_INCLUDE_DIRECTORIES or via the legacy variable +# `hiprtc_INCLUDE_DIR`. Pick whichever is populated; both end up at the +# same `/opt/rocm/include`. +if (DEFINED hiprtc_INCLUDE_DIR) + target_include_directories(rocmlir-tuning-driver SYSTEM PRIVATE + ${hiprtc_INCLUDE_DIR}) endif() +if (TARGET hiprtc::hiprtc) + get_target_property(_hiprtc_includes hiprtc::hiprtc + INTERFACE_INCLUDE_DIRECTORIES) + if (_hiprtc_includes) + target_include_directories(rocmlir-tuning-driver SYSTEM PRIVATE + ${_hiprtc_includes}) + endif() + unset(_hiprtc_includes) endif() +# Suppress HIP header warnings (HIP headers use GNU extensions that +# trigger -pedantic / -Wnested-anon-types / -Wreturn-type-c-linkage). +foreach(_src rocmlir-tuning-driver.cpp CacheFlush.cpp HipDelayLoad.cpp) + set_source_files_properties(${_src} PROPERTIES + COMPILE_OPTIONS + "-Wno-c++98-compat-extra-semi;-Wno-return-type-c-linkage;-Wno-nested-anon-types;-Wno-gnu-anonymous-struct") +endforeach() + set_property(TARGET rocmlir-tuning-driver PROPERTY INSTALL_RPATH_USE_LINK_PATH ON) +# `MLIRRocmRuntimeLoader` is the shared static library that owns the +# dlmopen / dlopen / LoadLibraryW machinery (and pulls `${CMAKE_DL_LIBS}` +# transitively). target_link_libraries(rocmlir-tuning-driver PRIVATE ${LIBS}) -target_link_libraries(rocmlir-tuning-driver PUBLIC benchmark-driver-utils) -target_link_libraries(rocmlir-tuning-driver PUBLIC hip::host hip::amdhip64 hiprtc::hiprtc) +target_link_libraries(rocmlir-tuning-driver PRIVATE MLIRRocmRuntimeLoader) llvm_update_compile_flags(rocmlir-tuning-driver) endif() diff --git a/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp b/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp index fd82c080ab5d..6c823ed09964 100644 --- a/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp +++ b/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp @@ -22,6 +22,12 @@ #include #endif +// Resolve every hipXXX/hiprtcXXX entry point through the delay-load +// function table so this TU never link-pulls libamdhip64 / libhiprtc. +// MUST be the last include block in this file (see header for rationale). +#include "HipDelayLoad.h" +#include "HipDelayLoadMacros.h" + using namespace mlir; namespace rocmlir::tuningdriver { diff --git a/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.cpp b/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.cpp new file mode 100644 index 000000000000..47c8dced484c --- /dev/null +++ b/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.cpp @@ -0,0 +1,177 @@ +//===- HipDelayLoad.cpp - Lazy HIP/HIPRTC symbol resolution ---------------===// +// +// Part of the rocMLIR 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 +// +//===----------------------------------------------------------------------===// + +#include "HipDelayLoad.h" + +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" + +#include +#include + +namespace rocmlir::tuningdriver { + +namespace { + +// Fail-fast helper. The tuning driver is fundamentally useless without HIP / +// HIPRTC -- every benchmark needs to launch a kernel on an AMD GPU. If we +// returned a partially-initialised symbol table here, the macros in +// `HipDelayLoadMacros.h` would later dereference a null function pointer and +// segfault. Aborting at first detection turns the failure mode into a clear +// diagnostic instead of an undebuggable crash. +[[noreturn]] void abortMissingHip(const char *what) { + std::fprintf( + stderr, + "rocmlir-tuning-driver: %s. The tuning driver requires a working " + "ROCm install; aborting.\n", + what); + std::abort(); +} + +HipSymbols loadHipSymbols() { + HipSymbols s; + s.lib = mlir::rocm_loader::loadRocmLibrary(mlir::rocm_loader::Library::Hip); + if (!s.lib.handle) + abortMissingHip("libamdhip64 not found on the loader search path " + "(tried unversioned alias and `.so.` for MAJOR " + "in [99..1])"); + + auto resolve = [&](const char *name) { + return mlir::rocm_loader::resolveRocmSymbol(s.lib, name); + }; + +#define LOAD_HIP_SYM(FIELD, NAME, TYPE) \ + do { \ + s.FIELD = reinterpret_cast(resolve(NAME)); \ + if (!s.FIELD) \ + abortMissingHip("missing required HIP symbol '" NAME \ + "' in libamdhip64"); \ + } while (false) + + LOAD_HIP_SYM(getDevice, "hipGetDevice", hipError_t (*)(int *)); + // hipGetDeviceProperties was renamed to ...R0600 for ABI stability in + // ROCm 6.0; prefer the new name and fall back to the legacy alias. + s.getDeviceProperties = + reinterpret_cast( + resolve("hipGetDevicePropertiesR0600")); + if (!s.getDeviceProperties) { + s.getDeviceProperties = + reinterpret_cast( + resolve("hipGetDeviceProperties")); + } + if (!s.getDeviceProperties) + abortMissingHip("neither 'hipGetDevicePropertiesR0600' nor " + "'hipGetDeviceProperties' found in libamdhip64"); + LOAD_HIP_SYM(getLastError, "hipGetLastError", hipError_t (*)(void)); + LOAD_HIP_SYM(getErrorString, "hipGetErrorString", + const char *(*)(hipError_t)); + + LOAD_HIP_SYM(malloc_, "hipMalloc", hipError_t (*)(void **, size_t)); + LOAD_HIP_SYM(free_, "hipFree", hipError_t (*)(void *)); + LOAD_HIP_SYM(memsetAsync, "hipMemsetAsync", + hipError_t (*)(void *, int, size_t, hipStream_t)); + + LOAD_HIP_SYM(streamCreate, "hipStreamCreate", hipError_t (*)(hipStream_t *)); + LOAD_HIP_SYM(streamDestroy, "hipStreamDestroy", hipError_t (*)(hipStream_t)); + LOAD_HIP_SYM(streamSynchronize, "hipStreamSynchronize", + hipError_t (*)(hipStream_t)); + + LOAD_HIP_SYM(eventCreate, "hipEventCreate", hipError_t (*)(hipEvent_t *)); + LOAD_HIP_SYM(eventDestroy, "hipEventDestroy", hipError_t (*)(hipEvent_t)); + LOAD_HIP_SYM(eventSynchronize, "hipEventSynchronize", + hipError_t (*)(hipEvent_t)); + LOAD_HIP_SYM(eventElapsedTime, "hipEventElapsedTime", + hipError_t (*)(float *, hipEvent_t, hipEvent_t)); + + LOAD_HIP_SYM(moduleLoadData, "hipModuleLoadData", + hipError_t (*)(hipModule_t *, const void *)); + LOAD_HIP_SYM(moduleUnload, "hipModuleUnload", hipError_t (*)(hipModule_t)); + LOAD_HIP_SYM(moduleGetFunction, "hipModuleGetFunction", + hipError_t (*)(hipFunction_t *, hipModule_t, const char *)); + LOAD_HIP_SYM(moduleLaunchKernel, "hipModuleLaunchKernel", + hipError_t (*)(hipFunction_t, unsigned, unsigned, unsigned, + unsigned, unsigned, unsigned, unsigned, + hipStream_t, void **, void **)); + LOAD_HIP_SYM(extModuleLaunchKernel, "hipExtModuleLaunchKernel", + hipError_t (*)(hipFunction_t, uint32_t, uint32_t, uint32_t, + uint32_t, uint32_t, uint32_t, size_t, hipStream_t, + void **, void **, hipEvent_t, hipEvent_t, + uint32_t)); +#undef LOAD_HIP_SYM + return s; +} + +#if defined(__HIP_PLATFORM_AMD__) +[[noreturn]] void abortMissingHiprtc(const char *what) { + std::fprintf(stderr, + "rocmlir-tuning-driver: %s. Runtime kernel compilation needs " + "libhiprtc; aborting.\n", + what); + std::abort(); +} + +HiprtcSymbols loadHiprtcSymbols(void *hipHandle) { + HiprtcSymbols s; + // HIPRTC shares HIP's KFD session (it only JIT-compiles GPU code; + // it does not open a separate device). Load it into HIP's link-map + // namespace so the same HSA instance satisfies both. + s.lib = mlir::rocm_loader::loadRocmLibrary(mlir::rocm_loader::Library::Hiprtc, + hipHandle); + if (!s.lib.handle) + abortMissingHiprtc("libhiprtc not found on the loader search path " + "(tried unversioned alias and `.so.` for " + "MAJOR in [99..1])"); + + auto resolve = [&](const char *name) { + return mlir::rocm_loader::resolveRocmSymbol(s.lib, name); + }; + +#define LOAD_HIPRTC_SYM(FIELD, NAME, TYPE) \ + do { \ + s.FIELD = reinterpret_cast(resolve(NAME)); \ + if (!s.FIELD) \ + abortMissingHiprtc("missing required HIPRTC symbol '" NAME \ + "' in libhiprtc"); \ + } while (false) + + LOAD_HIPRTC_SYM(getErrorString, "hiprtcGetErrorString", + const char *(*)(hiprtcResult)); + LOAD_HIPRTC_SYM(createProgram, "hiprtcCreateProgram", + hiprtcResult (*)(hiprtcProgram *, const char *, const char *, + int, const char **, const char **)); + LOAD_HIPRTC_SYM(destroyProgram, "hiprtcDestroyProgram", + hiprtcResult (*)(hiprtcProgram *)); + LOAD_HIPRTC_SYM(compileProgram, "hiprtcCompileProgram", + hiprtcResult (*)(hiprtcProgram, int, const char **)); + LOAD_HIPRTC_SYM(getProgramLogSize, "hiprtcGetProgramLogSize", + hiprtcResult (*)(hiprtcProgram, size_t *)); + LOAD_HIPRTC_SYM(getProgramLog, "hiprtcGetProgramLog", + hiprtcResult (*)(hiprtcProgram, char *)); + LOAD_HIPRTC_SYM(getCodeSize, "hiprtcGetCodeSize", + hiprtcResult (*)(hiprtcProgram, size_t *)); + LOAD_HIPRTC_SYM(getCode, "hiprtcGetCode", + hiprtcResult (*)(hiprtcProgram, char *)); +#undef LOAD_HIPRTC_SYM + return s; +} +#endif // __HIP_PLATFORM_AMD__ + +} // namespace + +const HipSymbols &getHipSymbols() { + static HipSymbols syms = loadHipSymbols(); + return syms; +} + +#if defined(__HIP_PLATFORM_AMD__) +const HiprtcSymbols &getHiprtcSymbols() { + static HiprtcSymbols syms = loadHiprtcSymbols(getHipSymbols().lib.handle); + return syms; +} +#endif + +} // namespace rocmlir::tuningdriver diff --git a/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.h b/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.h new file mode 100644 index 000000000000..650b02527c6b --- /dev/null +++ b/mlir/tools/rocmlir-tuning-driver/HipDelayLoad.h @@ -0,0 +1,143 @@ +//===- HipDelayLoad.h - Lazy HIP/HIPRTC symbol resolution -------*- C++ -*-===// +// +// Part of the rocMLIR 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 +// +//===----------------------------------------------------------------------===// +// +// Resolves the HIP runtime and HIPRTC entry points used by +// `rocmlir-tuning-driver` at run time rather than at link time. +// +// rocMLIR's executables embed their own LLVM (libLLVMSupport.so.*, +// libLLVMCodeGen.so.*, ...). Linking `libamdhip64` directly drags in +// `libamd_comgr` and ROCm's monolithic `libLLVM.so.`, which the +// dynamic linker maps into the host process at startup. When that LLVM +// runs its `cl::opt` static initializers, the dynamic linker unifies +// the rocMLIR-side and ROCm-side `cl::*` symbols across the split-vs- +// monolithic libraries. The result is a SmallPtrSet "Bucket < End" +// assertion (or `LLVM ERROR: Option '...' already exists!`) firing in +// `_dl_init`, before `main()` is even reached. See the dlopenHip +// branch's `[EXTERNAL] Coordinate HIP namespace ...` commit for the +// canonical write-up. +// +// Each consumer .cpp follows this pattern: +// +// #include +// #include // optional, for hipExtModuleLaunchKernel +// #include // optional, for HIPRTC users +// #include "HipDelayLoad.h" +// #include "HipDelayLoadMacros.h" // sequence of #define hipXXX(...) ... +// +// after which bare `hipMalloc(...)`, `hiprtcCreateProgram(...)` etc. +// expand to the function-pointer table dispatch; HIP / HIPRTC types +// (`hipModule_t`, `hipDeviceProp_t`, `hiprtcResult`, ...) come from the +// real headers and are unchanged. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_TOOLS_ROCMLIR_TUNING_DRIVER_HIPDELAYLOAD_H +#define MLIR_TOOLS_ROCMLIR_TUNING_DRIVER_HIPDELAYLOAD_H + +// We need the full HIP / HIPRTC type definitions to express the +// function-pointer signatures below (`hipModule_t`, `hipDeviceProp_t`, +// `hiprtcProgram`, ...). Pull them in directly so consumers can include +// `HipDelayLoad.h` unconditionally without first including the HIP +// headers themselves. HIP headers are header-only for type purposes +// (no transitive linkage); the dlmopen-based loader in HipDelayLoad.cpp +// is what keeps libamdhip64 out of the build-time link line. +#include +#include +#if defined(__HIP_PLATFORM_AMD__) +#include +#endif + +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" + +#include +#include + +namespace rocmlir::tuningdriver { + +/// Function pointers for every HIP entry point used by the tuning driver. +/// `lib.handle` is non-null and every function pointer is non-null on a +/// successful load -- the loader aborts the process via +/// `std::abort()` rather than returning a partially-populated table, +/// because the bare-call macros in `HipDelayLoadMacros.h` would otherwise +/// dispatch through a null function pointer (UB / segfault) at the call +/// site. The tuning driver is fundamentally useless without HIP. +struct HipSymbols { + mlir::rocm_loader::LoadedLibrary lib; + + hipError_t (*getDevice)(int *) = nullptr; + hipError_t (*getDeviceProperties)(hipDeviceProp_t *, int) = nullptr; + hipError_t (*getLastError)(void) = nullptr; + const char *(*getErrorString)(hipError_t) = nullptr; + + hipError_t (*malloc_)(void **, size_t) = nullptr; + hipError_t (*free_)(void *) = nullptr; + hipError_t (*memsetAsync)(void *, int, size_t, hipStream_t) = nullptr; + + hipError_t (*streamCreate)(hipStream_t *) = nullptr; + hipError_t (*streamDestroy)(hipStream_t) = nullptr; + hipError_t (*streamSynchronize)(hipStream_t) = nullptr; + + hipError_t (*eventCreate)(hipEvent_t *) = nullptr; + hipError_t (*eventDestroy)(hipEvent_t) = nullptr; + hipError_t (*eventSynchronize)(hipEvent_t) = nullptr; + hipError_t (*eventElapsedTime)(float *, hipEvent_t, hipEvent_t) = nullptr; + + hipError_t (*moduleLoadData)(hipModule_t *, const void *) = nullptr; + hipError_t (*moduleUnload)(hipModule_t) = nullptr; + hipError_t (*moduleGetFunction)(hipFunction_t *, hipModule_t, + const char *) = nullptr; + hipError_t (*moduleLaunchKernel)(hipFunction_t, unsigned, unsigned, unsigned, + unsigned, unsigned, unsigned, unsigned, + hipStream_t, void **, void **) = nullptr; + hipError_t (*extModuleLaunchKernel)(hipFunction_t, uint32_t, uint32_t, + uint32_t, uint32_t, uint32_t, uint32_t, + size_t, hipStream_t, void **, void **, + hipEvent_t, hipEvent_t, + uint32_t) = nullptr; +}; + +/// Process-wide accessor for the HIP function table. Initialised on first +/// call. The implementation tries to reuse the dlmopen handle owned by +/// `RocmSystemDetect` (via `mlirRocmSystemDetectGetHipHandle` looked up +/// through `RTLD_DEFAULT`) so that we share a single HSA session per +/// process; KFD only permits one. If that symbol is absent (binary built +/// without RocmSystemDetect), falls back to its own dlmopen. The loader +/// `std::abort()`s the process on failure -- it never returns a +/// partially-populated table -- because the bare-call macros in +/// `HipDelayLoadMacros.h` would dispatch through a null function +/// pointer at the call site otherwise. +const HipSymbols &getHipSymbols(); + +#if defined(__HIP_PLATFORM_AMD__) +/// Same fail-fast contract as `HipSymbols`: every member is non-null on a +/// successful load; the loader aborts on missing libhiprtc or missing +/// required symbol. +struct HiprtcSymbols { + mlir::rocm_loader::LoadedLibrary lib; + + const char *(*getErrorString)(hiprtcResult) = nullptr; + hiprtcResult (*createProgram)(hiprtcProgram *, const char *, const char *, + int, const char **, const char **) = nullptr; + hiprtcResult (*destroyProgram)(hiprtcProgram *) = nullptr; + hiprtcResult (*compileProgram)(hiprtcProgram, int, const char **) = nullptr; + hiprtcResult (*getProgramLogSize)(hiprtcProgram, size_t *) = nullptr; + hiprtcResult (*getProgramLog)(hiprtcProgram, char *) = nullptr; + hiprtcResult (*getCodeSize)(hiprtcProgram, size_t *) = nullptr; + hiprtcResult (*getCode)(hiprtcProgram, char *) = nullptr; +}; + +/// Process-wide accessor for the HIPRTC function table. HIPRTC ships in +/// its own SONAME (`libhiprtc.so.`), so it is loaded via a +/// separate dlmopen / dlopen call from HIP. HIPRTC is only used by the +/// instruction-cache flush kernel JIT in CacheFlush.cpp. +const HiprtcSymbols &getHiprtcSymbols(); +#endif + +} // namespace rocmlir::tuningdriver + +#endif // MLIR_TOOLS_ROCMLIR_TUNING_DRIVER_HIPDELAYLOAD_H diff --git a/mlir/tools/rocmlir-tuning-driver/HipDelayLoadMacros.h b/mlir/tools/rocmlir-tuning-driver/HipDelayLoadMacros.h new file mode 100644 index 000000000000..13ecece5c27b --- /dev/null +++ b/mlir/tools/rocmlir-tuning-driver/HipDelayLoadMacros.h @@ -0,0 +1,104 @@ +//===- HipDelayLoadMacros.h - Macros redirecting hipXXX/hiprtcXXX -*- C++-*-=// +// +// Part of the rocMLIR 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 +// +//===----------------------------------------------------------------------===// +// +// Bag of preprocessor `#define`s that turn bare `hipXXX(args)` / +// `hiprtcXXX(args)` call sites into dispatches through the function- +// pointer table provided by `HipDelayLoad.h`. The macros expand to the +// canonical C-call form; HIP's own header definitions are simply +// shadowed. +// +// Include this header LAST in any TU that wants to use the redirects. +// In particular, do NOT include any HIP header *after* this file -- +// doing so would either redefine the wrapped function symbols (link +// error) or, worse, silently restore the direct linkage. The expected +// pattern is: +// +// #include +// #include +// #include +// #include "HipDelayLoad.h" +// #include "HipDelayLoadMacros.h" +// // ... rest of TU uses bare hipMalloc(...), hiprtcCompileProgram(...) ... +// +// This file is intentionally not include-guarded: every consumer that +// includes it gets a fresh expansion at the point of inclusion. Re- +// including it within the same TU would just redefine the same macros +// to the same expansions and is harmless (modulo a -Wmacro-redefined +// warning). +// +//===----------------------------------------------------------------------===// + +// HIP runtime API +#define hipGetDevice(...) \ + (::rocmlir::tuningdriver::getHipSymbols().getDevice(__VA_ARGS__)) +#define hipGetDeviceProperties(...) \ + (::rocmlir::tuningdriver::getHipSymbols().getDeviceProperties(__VA_ARGS__)) +#define hipGetLastError(...) \ + (::rocmlir::tuningdriver::getHipSymbols().getLastError(__VA_ARGS__)) +#define hipGetErrorString(...) \ + (::rocmlir::tuningdriver::getHipSymbols().getErrorString(__VA_ARGS__)) + +#define hipMalloc(...) \ + (::rocmlir::tuningdriver::getHipSymbols().malloc_(__VA_ARGS__)) +#define hipFree(...) \ + (::rocmlir::tuningdriver::getHipSymbols().free_(__VA_ARGS__)) +#define hipMemsetAsync(...) \ + (::rocmlir::tuningdriver::getHipSymbols().memsetAsync(__VA_ARGS__)) + +#define hipStreamCreate(...) \ + (::rocmlir::tuningdriver::getHipSymbols().streamCreate(__VA_ARGS__)) +#define hipStreamDestroy(...) \ + (::rocmlir::tuningdriver::getHipSymbols().streamDestroy(__VA_ARGS__)) +#define hipStreamSynchronize(...) \ + (::rocmlir::tuningdriver::getHipSymbols().streamSynchronize(__VA_ARGS__)) + +#define hipEventCreate(...) \ + (::rocmlir::tuningdriver::getHipSymbols().eventCreate(__VA_ARGS__)) +#define hipEventDestroy(...) \ + (::rocmlir::tuningdriver::getHipSymbols().eventDestroy(__VA_ARGS__)) +#define hipEventSynchronize(...) \ + (::rocmlir::tuningdriver::getHipSymbols().eventSynchronize(__VA_ARGS__)) +#define hipEventElapsedTime(...) \ + (::rocmlir::tuningdriver::getHipSymbols().eventElapsedTime(__VA_ARGS__)) + +#define hipModuleLoadData(...) \ + (::rocmlir::tuningdriver::getHipSymbols().moduleLoadData(__VA_ARGS__)) +#define hipModuleUnload(...) \ + (::rocmlir::tuningdriver::getHipSymbols().moduleUnload(__VA_ARGS__)) +#define hipModuleGetFunction(...) \ + (::rocmlir::tuningdriver::getHipSymbols().moduleGetFunction(__VA_ARGS__)) +#define hipModuleLaunchKernel(...) \ + (::rocmlir::tuningdriver::getHipSymbols().moduleLaunchKernel(__VA_ARGS__)) +// The real `hipExtModuleLaunchKernel` takes 14 args; the last (`flags`) +// defaults to 0 in the HIP header. Default arguments do not apply when +// calling through a function pointer, so existing 13-arg call sites +// would fail to type-check. Keep the call-site syntax identical to +// upstream HIP code by injecting the trailing `0u` here. +#define hipExtModuleLaunchKernel(...) \ + (::rocmlir::tuningdriver::getHipSymbols().extModuleLaunchKernel( \ + __VA_ARGS__, 0u)) + +// HIPRTC API (only on AMD; HIPRTC has no NVIDIA-side stub). +#if defined(__HIP_PLATFORM_AMD__) +#define hiprtcGetErrorString(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().getErrorString(__VA_ARGS__)) +#define hiprtcCreateProgram(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().createProgram(__VA_ARGS__)) +#define hiprtcDestroyProgram(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().destroyProgram(__VA_ARGS__)) +#define hiprtcCompileProgram(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().compileProgram(__VA_ARGS__)) +#define hiprtcGetProgramLogSize(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().getProgramLogSize(__VA_ARGS__)) +#define hiprtcGetProgramLog(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().getProgramLog(__VA_ARGS__)) +#define hiprtcGetCodeSize(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().getCodeSize(__VA_ARGS__)) +#define hiprtcGetCode(...) \ + (::rocmlir::tuningdriver::getHiprtcSymbols().getCode(__VA_ARGS__)) +#endif diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index b2794d9c6881..a7ec2848e0b8 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -85,6 +85,12 @@ void pArgs(const std::tuple &formals, void **_vargs) { // Needs to go second lest we get compiler issues #include +// Resolve every hipXXX entry point through the delay-load function +// table so this TU never link-pulls libamdhip64. MUST be the last +// include block (see HipDelayLoad.h for rationale). +#include "HipDelayLoad.h" +#include "HipDelayLoadMacros.h" + using namespace mlir; using namespace rocmlir::tuningdriver; diff --git a/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp b/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp index c30ac8f152b4..c35e37d9cfde 100644 --- a/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp +++ b/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp @@ -10,38 +10,47 @@ #include "gtest/gtest.h" -#include "hip/hip_runtime_api.h" - #include +#include +#include using namespace mlir::rock; -class NativeArchTest : public ::testing::TestWithParam { +// NOTE: this file deliberately does NOT include hip/hip_runtime_api.h or link +// libamdhip64. Doing so would pull in amd_comgr -> ROCm's libLLVM.so, which +// collides with the LLVM that the test binary embeds. Device enumeration and +// arch-name lookup go through the AmdArchDb public API, which delay-loads HIP +// at run time via a private `dlopen` / `LoadLibraryW` call. + +class NativeArchTest : public ::testing::TestWithParam { public: static auto getDeviceIds() { - int count; - if (auto err = hipGetDeviceCount(&count); err != hipSuccess) { - return ::testing::ValuesIn({0}); + unsigned count = nativeDeviceCount(); + if (count == 0) { + // Keep gtest happy when no GPU/HIP is available; the SetUp() below will + // skip the test for the synthetic device id. + return ::testing::ValuesIn(std::vector{0}); } - std::vector ids(count); - std::iota(ids.begin(), ids.end(), 0); + std::vector ids(count); + std::iota(ids.begin(), ids.end(), 0u); return ::testing::ValuesIn(ids); } protected: void SetUp() override { - if (auto err = hipGetDeviceProperties(&prop, GetParam()); - err != hipSuccess) { - FAIL() << "hipGetDeviceProperties failed with error: " - << hipGetErrorString(err); - } + archName = nativeArchName(GetParam()); + if (archName.empty()) + GTEST_SKIP() << "No AMD GPU visible to HIP (or `libamdhip64` not on the " + "loader path); skipping native arch comparison for " + "device " + << GetParam(); } - hipDeviceProp_t prop; + std::string archName; }; TEST_P(NativeArchTest, NativeArchInfoMatchesPresetInfo) { - auto presetInfo = lookupArchInfo(prop.gcnArchName); + auto presetInfo = lookupArchInfo(archName); auto nativeInfo = lookupArchInfo("native:" + std::to_string(GetParam())); EXPECT_EQ(presetInfo.defaultFeatures, nativeInfo.defaultFeatures); @@ -62,3 +71,85 @@ TEST_P(NativeArchTest, NativeArchInfoMatchesPresetInfo) { INSTANTIATE_TEST_SUITE_P(NativeArchTests, NativeArchTest, NativeArchTest::getDeviceIds()); + +// Pin the parser contract for `rock.arch = "native[:N]"`. Malformed input +// (`native:foo`, `native:`, `native:-1`, `native:9999999999999999999999`) used +// to silently fall back to device 0, which on multi-GPU systems silently +// targeted the wrong GPU. The parser must abort instead. +// +// We use `EXPECT_DEATH` so the test works whether or not a real GPU is +// available -- the abort happens before any HIP call. +TEST(NativeArchParseTest, MalformedSuffixAborts) { + EXPECT_DEATH( + { (void)lookupArchInfo("native:foo"); }, + "Invalid `rock.arch = \"native:foo\"`"); + EXPECT_DEATH( + { (void)lookupArchInfo("native:1abc"); }, + "Invalid `rock.arch = \"native:1abc\"`"); + EXPECT_DEATH( + { (void)lookupArchInfo("native:"); }, + "Invalid `rock.arch = \"native:\"`"); + EXPECT_DEATH( + { (void)lookupArchInfo("native:-1"); }, + "Invalid `rock.arch = \"native:-1\"`"); +} + +// Bare `native` (no colon) is well-formed and means "device 0". +TEST(NativeArchParseTest, BareNativeIsDeviceZero) { + // We cannot directly observe the parsed deviceId without a GPU, but we can + // at least confirm the parse does not abort. If HIP is unavailable, the + // call later aborts with a *different* message ("Failed to query AMD GPU + // arch runtime"), which is still a valid outcome distinct from the parser + // abort above. + if (nativeDeviceCount() == 0) + GTEST_SKIP() << "no AMD GPU visible; the parse-success path needs a " + "live HIP runtime to return without aborting"; + // Should not abort. + (void)lookupArchInfo("native"); +} + +// On multi-GPU systems with same-arch devices, the per-device cache must +// not collapse data across device ids. The previous implementation keyed +// the cache by `gcnArchName` only and silently returned device 0's data +// for every later device. +// +// Skipped when fewer than two visible GPUs share the same arch name (which +// is the common single-GPU CI case). +TEST(NativeArchCacheTest, SameArchMultiGpuDistinct) { + unsigned count = nativeDeviceCount(); + if (count < 2) + GTEST_SKIP() << "fewer than 2 AMD GPUs visible; cannot exercise the " + "same-arch multi-GPU cache contract"; + + std::string arch0 = nativeArchName(0); + if (arch0.empty()) + GTEST_SKIP() << "device 0 unavailable; cannot exercise the cache"; + + // Find a second device that reports the same gcnArchName as device 0. + unsigned other = count; + for (unsigned i = 1; i < count; ++i) { + if (nativeArchName(i) == arch0) { + other = i; + break; + } + } + if (other == count) + GTEST_SKIP() << "no two visible GPUs share `" << arch0 << "`"; + + auto info0 = lookupArchInfo("native:0"); + auto infoN = lookupArchInfo("native:" + std::to_string(other)); + + // The per-device CU count is the canonical "is the cache device-aware?" + // probe: even on otherwise-identical SKUs, AMD's binning can produce + // different `multiProcessorCount` (= minNumCU after our query). Two + // CALLS to `lookupArchInfo` for two distinct device ids must land on + // independently queried entries, not on a stale device-0 copy. + // + // We don't EXPECT_NE here because two physically identical GPUs may also + // legitimately return the same minNumCU; the meaningful invariant is that + // each value comes from its own per-device query and is consistent on + // repeated lookup. Repeat the call to prove cache stability: + EXPECT_EQ(lookupArchInfo("native:0").minNumCU, info0.minNumCU); + EXPECT_EQ(lookupArchInfo("native:" + std::to_string(other)).minNumCU, + infoN.minNumCU); +} diff --git a/mlir/unittests/Dialect/Rock/CMakeLists.txt b/mlir/unittests/Dialect/Rock/CMakeLists.txt index 5f2f03d432e7..ab615968d39a 100644 --- a/mlir/unittests/Dialect/Rock/CMakeLists.txt +++ b/mlir/unittests/Dialect/Rock/CMakeLists.txt @@ -1,4 +1,5 @@ set(ROCK_UNITTEST_SOURCES + AmdArchDbTests.cpp TransformMapBuilderTests.cpp TosaUtilsTests.cpp loweringUtilsTests.cpp @@ -9,8 +10,16 @@ set(ROCK_UNITTEST_SOURCES ParamLookupTableTests.cpp ) -if(NOT WIN32) - list(APPEND ROCK_UNITTEST_SOURCES AmdArchDbTests.cpp) +# `RocmRuntimeLoaderTests` exercises the cross-library coordination +# contract between `MLIRRocmRuntimeLoader` and `RocmSystemDetect`. The +# latter lives in `MLIRRocmExecutionEngineUtils`, which is only built +# as a standalone target in the standard shared-library configuration. +# In the fat-library configuration (`BUILD_FAT_LIBROCKCOMPILER=ON`), +# `RocmSystemDetect` is not bundled into `librockCompiler.a` -- it is +# only exposed to the rocMLIR JIT path -- so we omit the test sub-suite +# there. The contract being tested is identical in both configurations. +if (TARGET MLIRRocmExecutionEngineUtils) + list(APPEND ROCK_UNITTEST_SOURCES RocmRuntimeLoaderTests.cpp) endif() add_rocmlir_unittest(MLIRRockUnitTests @@ -26,4 +35,8 @@ target_link_libraries(MLIRRockUnitTests MLIRTosaDialect MLIRFuncDialect MLIRIR + MLIRRocmRuntimeLoader ) +if (TARGET MLIRRocmExecutionEngineUtils) + target_link_libraries(MLIRRockUnitTests PRIVATE MLIRRocmExecutionEngineUtils) +endif() diff --git a/mlir/unittests/Dialect/Rock/RocmRuntimeLoaderTests.cpp b/mlir/unittests/Dialect/Rock/RocmRuntimeLoaderTests.cpp new file mode 100644 index 000000000000..07ef31d387ec --- /dev/null +++ b/mlir/unittests/Dialect/Rock/RocmRuntimeLoaderTests.cpp @@ -0,0 +1,128 @@ +//===- RocmRuntimeLoaderTests.cpp - tests for the ROCm runtime loader ----===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// These tests document the cross-process coordination contract of +// `mlir::rocm_loader::loadRocmLibrary`: any second consumer of HIP in +// the same process must observe the same handle as `RocmSystemDetect`, +// otherwise KFD's "one HSA session per process" rule produces +// `hipErrorNoDevice` for one of them. +// +// The tests are runtime-conditional: when no HIP runtime is installed +// (CI without `libamdhip64.so` on the loader path), they GTEST_SKIP +// instead of failing. +// +//===----------------------------------------------------------------------===// + +#include "mlir/ExecutionEngine/RocmRuntimeLoader.h" +#include "mlir/ExecutionEngine/RocmSystemDetect.h" + +#include "gtest/gtest.h" + +using mlir::rocm_loader::CoordinationPolicy; +using mlir::rocm_loader::Library; +using mlir::rocm_loader::LoadedLibrary; +using mlir::rocm_loader::loadRocmLibrary; +using mlir::rocm_loader::resolveRocmSymbol; + +namespace { + +// Touching `RocmSystemDetect::get()` forces it (the canonical owner) to +// load HIP in its own link-map namespace and publishes +// `mlirRocmSystemDetectGetHipHandle` for subsequent loaders to find. +// +// We do this in a SetUp helper rather than in the constructor so the +// test can GTEST_SKIP cleanly when HIP is unavailable. +class RocmRuntimeLoaderTest : public ::testing::Test { +protected: + void SetUp() override { + (void)mlir::RocmSystemDetect::get(); + sharedHandle = mlirRocmSystemDetectGetHipHandle(); + if (!sharedHandle) + GTEST_SKIP() << "no HIP runtime available; skipping loader contract " + "tests"; + } + + void *sharedHandle = nullptr; +}; + +// `CoordinationPolicy::Auto` for `Library::Hip` MUST return the same +// handle as `RocmSystemDetect`. If a future change accidentally opens a +// second HIP namespace via `dlmopen`, KFD will start handing out +// `hipErrorNoDevice` and large parts of the JIT path will break. +TEST_F(RocmRuntimeLoaderTest, AutoPolicyReusesSystemDetectHandle) { + LoadedLibrary lib = loadRocmLibrary(Library::Hip); + ASSERT_NE(nullptr, lib.handle) + << "loadRocmLibrary(Hip) returned a null handle even though " + "RocmSystemDetect successfully loaded HIP"; + EXPECT_EQ(sharedHandle, lib.handle) + << "Auto policy must reuse RocmSystemDetect's HIP handle to keep " + "the per-process KFD session count at 1"; +} + +// `CoordinationPolicy::Owned` is reserved for `RocmSystemDetect` itself +// (to break recursion at first load). When called from anywhere else +// it MAY return a fresh handle, but MUST still produce a usable one. +TEST_F(RocmRuntimeLoaderTest, OwnedPolicyAlwaysReturnsAUsableHandle) { + LoadedLibrary lib = loadRocmLibrary(Library::Hip, /*relatedHandle=*/nullptr, + CoordinationPolicy::Owned); + ASSERT_NE(nullptr, lib.handle) + << "loadRocmLibrary(Hip, Owned) failed even though HIP is present"; + // We do NOT assert handle equality here: under glibc, a fresh + // dlmopen call produces a distinct handle by design. The test only + // proves that Owned does not regress to a null result. + void *sym = resolveRocmSymbol(lib, "hipGetDeviceCount"); + EXPECT_NE(nullptr, sym) + << "An owned HIP handle must be able to resolve hipGetDeviceCount"; +} + +// Symbol resolution against a null handle must be a soft failure +// (return null), never a crash. This protects all the `if (!fn) return +// false;` fallbacks in the wrapper translation units. +TEST(RocmRuntimeLoaderUnit, ResolveAgainstNullHandleReturnsNullSafely) { + LoadedLibrary lib; + EXPECT_EQ(nullptr, resolveRocmSymbol(lib, "anything")); +} + +// Pin the version-agnosticism contract: if HIP is available at all on +// the host, the loader MUST find it without any compile-time knowledge +// of which ROCm major version is installed. We use the `Auto` policy +// (the path downstream consumers actually take) so we avoid bumping +// into KFD's "one HSA session per process" limit, which would mask a +// true loader failure under a spurious `dlmopen` failure from +// repeated `Owned` calls. Success here means the loader resolved a +// SONAME -- either via the cross-library coordination handle or via +// the bare-name / numeric-fallback enumeration. A future ROCm release +// that bumps the major version should keep this test green without +// code changes (so long as AMD stays at or below +// `kMaxProbedRocmMajor`). +class RocmRuntimeLoaderVersionContract : public ::testing::Test { +protected: + void SetUp() override { + (void)mlir::RocmSystemDetect::get(); + if (!mlirRocmSystemDetectGetHipHandle()) + GTEST_SKIP() << "no HIP runtime available on this host"; + } +}; + +TEST_F(RocmRuntimeLoaderVersionContract, FindsHipWithoutVersionHardcoding) { + LoadedLibrary lib = loadRocmLibrary(Library::Hip); + ASSERT_NE(nullptr, lib.handle) + << "loader could not resolve HIP through its candidate list " + "(bare name + numeric `.so.` fallback). Either ROCm is " + "missing from the dynamic-loader path, or AMD has shipped a " + "major version above `kMaxProbedRocmMajor` (in which case " + "bump that constant in RocmRuntimeLoader.cpp)"; + // `hipGetDeviceCount` has been part of the HIP C ABI since ROCm 1.x + // and is therefore present on every supported major version, which + // makes it the right symbol to pin "the loaded library is in fact + // HIP, not some other library that happened to match the SONAME + // pattern". + EXPECT_NE(nullptr, resolveRocmSymbol(lib, "hipGetDeviceCount")); +} + +} // namespace diff --git a/mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl b/mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl index 57ba9b62a598..003cfffbf73c 100755 --- a/mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl +++ b/mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl @@ -8,6 +8,18 @@ # It should be run from the build directory on Linux. +# Libraries that must NOT be bundled into librockCompiler.a, even if they +# appear in the dependency graph of MLIRRockThin. They are shipped as +# separate shared libraries on purpose -- bundling them would re-introduce +# the very dependencies the fat-lib split is meant to keep out. +# +# conv-validation-wrappers: a runner-only helper not consumed by +# downstream compile-time consumers (MIOpen, MIGraphX). Excluded from +# all-static fat-build mode via EXCLUDE_FROM_ALL ${BUILD_FAT_LIBROCKCOMPILER}. +my %excludedLibs = map { $_ => 1 } qw( + conv-validation-wrappers +); + my @rocmlirLibs; my @mlirLibs; @@ -17,9 +29,9 @@ foreach (@deps) { last if /outputs:/; if (m#external/llvm-project/llvm/lib/lib(\w+)\.a#) { - push @mlirLibs, $1; + push @mlirLibs, $1 unless $excludedLibs{$1}; } elsif (m#lib/lib(\w+)\.a#) { - push @rocmlirLibs, $1; + push @rocmlirLibs, $1 unless $excludedLibs{$1}; } } @mlirLibs = sort @mlirLibs;