From 23ad884acd9757b779ad92e3564563a8f02b9487 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sun, 10 May 2026 10:39:10 -0400 Subject: [PATCH] [Comgr] Skip embedded libc++ -idirafter when system C++ headers exist; fix VFS path mismatch with clang resource dir MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This commit fixes two related issues with comgr's embedded libc++ headers support that surface together when compiling HIP with `-nogpuinc` on real-world Linux hosts. Issue 1: libstdc++/libc++ header conflict (ROCm-issue-2445) When `-nogpuinc` is active and embedded libc++ headers are loaded into the VFS, comgr injected `-idirafter /include/c++/v1`. On distros that ship libstdc++ at clang's default include path (RHEL/manylinux/ Ubuntu), this produces a hybrid include chain. The host libstdc++ ``/`` pull in a transitive `` from the embedded libc++ overlay, whose `#include_next` then has no successor under `-nogpuinc` and fails. Example: `#include ` → libstdc++ `` → libc++ `stddef.h` (via VFS) → `#include_next ` → no next file → error. The embedded set is partial by design (LIBCXX_USER_HEADERS in cmake/LibcxxHeaders.cmake), so it can only safely substitute when no host C++ standard library is present. Fix: detect host libstdc++ / libc++ on disk (honoring `--sysroot=`) and skip the `-idirafter` injection when found. Also skip when the user passed `-nostdinc++`, `-nostdinc`, or `-nostdlibinc`. Add `AMD_COMGR_USE_EMBEDDED_LIBCXX={auto,force,disable}` env override. Issue 2: VFS embed paths drift from clang's resource-dir path The clang Driver and the VFS embed code each constructed a clang binary path from `LLVM_PATH`, but using different string-joining functions. When `LLVM_PATH` was unset, the two disagreed: Driver: `(Twine("") + "/bin/clang").str()` → "/bin/clang" Embed: `SmallString<256>(""); append("bin","clang")` → "bin/clang" Both then fed `clang::GetResourcesPath()`, which strips two parents and appends `lib/clang/`: "/bin/clang" → "/lib/clang/23" (absolute) "bin/clang" → "lib/clang/23" (relative) So clang searched `/lib/clang/23/include/stddef.h` while comgr planted the file at `lib/clang/23/include/stddef.h`. Mismatch → header not found. Fix: extract a single `env::getClangBinaryPath()` helper used by both the Driver and the VFS embed code, so the two paths cannot drift again. When `LLVM_PATH` is unset, locate the real clang via `dladdr` on the loaded libamd_comgr.so and probe sibling layouts (`/llvm/bin/ clang` for ROCm, `/bin/clang` for standard installs); fall back to `/bin/clang` only if neither is found, which keeps Driver and VFS consistent even on unusual layouts. Tests: added compile_hip_stdlibcxx_conflict_test (system-libstdc++ chain), compile_hip_with_system_libcxx_test (asserts skip-detection fires when system libc++ present), compile_hip_distroless_test (forces embedded path via env override). All three skip cleanly on Windows; the conflict test also skips when no clang resource dir is reachable on disk (no false failure when prerequisites aren't present). --- amd/comgr/src/comgr-compiler.cpp | 152 +++++++++-- amd/comgr/src/comgr-compiler.h | 9 + amd/comgr/src/comgr-env.cpp | 71 +++++ amd/comgr/src/comgr-env.h | 18 ++ amd/comgr/test/CMakeLists.txt | 3 + amd/comgr/test/compile_hip_distroless_test.c | 108 ++++++++ .../compile_hip_stdlibcxx_conflict_test.c | 252 ++++++++++++++++++ .../compile_hip_with_system_libcxx_test.c | 157 +++++++++++ 8 files changed, 754 insertions(+), 16 deletions(-) create mode 100644 amd/comgr/test/compile_hip_distroless_test.c create mode 100644 amd/comgr/test/compile_hip_stdlibcxx_conflict_test.c create mode 100644 amd/comgr/test/compile_hip_with_system_libcxx_test.c diff --git a/amd/comgr/src/comgr-compiler.cpp b/amd/comgr/src/comgr-compiler.cpp index 1b96d5ff908e3..c889cd1942d21 100644 --- a/amd/comgr/src/comgr-compiler.cpp +++ b/amd/comgr/src/comgr-compiler.cpp @@ -898,7 +898,7 @@ AMDGPUCompiler::executeInProcessDriver(ArrayRef Args) { ProcessWarningOptions(Diags, *DiagOpts, *OverlayFS, /*ReportDiags=*/false); - Driver TheDriver((Twine(env::getLLVMPath()) + "/bin/clang").str(), + Driver TheDriver(env::getClangBinaryPath(), llvm::sys::getDefaultTargetTriple(), Diags, "AMDGPU Code Object Manager", OverlayFS); TheDriver.setCheckInputsExist(false); @@ -1031,6 +1031,127 @@ amd_comgr_status_t AMDGPUCompiler::removeTmpDirs() { #endif } +// Probe a single directory tree for system C++ headers: +// /include/c++/v1/__config_site (libc++) +// /include/c++//cstddef (libstdc++) +// On hit, writes the matching path to `FoundPath` (if non-null) and returns +// true. Does not model Debian multiarch `g++-multiarch-incdir` layout, but +// `cstddef` itself lives in the common version dir on all major distros, so +// the libstdc++ probe still triggers there. +static bool probeCxxHeadersUnder(StringRef Root, std::string *FoundPath) { + auto Hit = [&](const Twine &P) { + if (FoundPath) + *FoundPath = P.str(); + return true; + }; + + SmallString<256> LibCxx(Root); + sys::path::append(LibCxx, "include", "c++", "v1", "__config_site"); + if (sys::fs::exists(LibCxx)) + return Hit(LibCxx); + + SmallString<256> CxxRoot(Root); + sys::path::append(CxxRoot, "include", "c++"); + std::error_code EC; + for (sys::fs::directory_iterator DI(CxxRoot, EC), End; DI != End && !EC; + DI.increment(EC)) { + if (DI->type() != sys::fs::file_type::directory_file) + continue; + SmallString<256> Probe(DI->path()); + sys::path::append(Probe, "cstddef"); + if (sys::fs::exists(Probe)) + return Hit(Probe); + } + return false; +} + +// Filesystem probe for system libstdc++ / libc++ headers. Honors user args +// that redirect clang's header search: +// --sysroot= → probe /usr{,/local}/... +// --gcc-toolchain= → probe /... +// Without a full Generic_GCC search-path model this misses exotic layouts +// (Gentoo crossdev, multilib biarch tuples, --gcc-install-dir relative +// climbs); in those cases we fall through to "no system headers detected" +// and inject embedded — same behavior as before this fix. +// +// On hit, `FoundPath` (if non-null) receives the matching header path for +// diagnostic logging. +static bool detectSystemCxxHeadersOnDisk(ArrayRef Argv, + std::string *FoundPath) { + std::string SysRoot; + std::string GccToolchain; + for (size_t I = 0; I < Argv.size(); ++I) { + StringRef A(Argv[I] ? Argv[I] : ""); + if (A == "--sysroot" && I + 1 < Argv.size() && Argv[I + 1]) + SysRoot = Argv[I + 1]; + else if (A.starts_with("--sysroot=")) + SysRoot = A.drop_front(StringRef("--sysroot=").size()).str(); + else if (A == "--gcc-toolchain" && I + 1 < Argv.size() && Argv[I + 1]) + GccToolchain = Argv[I + 1]; + else if (A.starts_with("--gcc-toolchain=")) + GccToolchain = A.drop_front(StringRef("--gcc-toolchain=").size()).str(); + } + if (SysRoot.empty()) + SysRoot = "/"; + + // GCC toolchain wins if specified — that's what clang's driver would + // resolve C++ headers under. + if (!GccToolchain.empty() && probeCxxHeadersUnder(GccToolchain, FoundPath)) + return true; + + SmallString<256> SysUsr(SysRoot); + sys::path::append(SysUsr, "usr"); + if (probeCxxHeadersUnder(SysUsr, FoundPath)) + return true; + + SmallString<256> SysUsrLocal(SysRoot); + sys::path::append(SysUsrLocal, "usr", "local"); + if (probeCxxHeadersUnder(SysUsrLocal, FoundPath)) + return true; + + return false; +} + +bool AMDGPUCompiler::shouldSkipEmbeddedHeaders(ArrayRef Argv) { + if (SkipEmbeddedHeadersCache) + return *SkipEmbeddedHeadersCache; + + bool Verbose = env::shouldEmitVerboseLogs(); + auto Decide = [&](bool Skip, const Twine &Reason) { + SkipEmbeddedHeadersCache = Skip; + if (Verbose) + LogS << "\t Embedded libc++ headers: " << (Skip ? "skipped" : "active") + << " (" << Reason << ")\n"; + return Skip; + }; + + // Env override takes precedence. + switch (env::getEmbeddedLibcxxMode()) { + case env::EmbeddedLibcxxMode::Force: + return Decide(false, "AMD_COMGR_USE_EMBEDDED_LIBCXX=force"); + case env::EmbeddedLibcxxMode::Disable: + return Decide(true, "AMD_COMGR_USE_EMBEDDED_LIBCXX=disable"); + case env::EmbeddedLibcxxMode::Auto: + break; + } + + // User explicitly took control of C++ include search — don't second-guess. + for (const char *A : Argv) { + if (!A) + continue; + StringRef S(A); + if (S == "-nostdinc++" || S == "-nostdinc" || S == "-nostdlibinc") + return Decide(true, Twine("user passed ") + S); + } + + // System C++ headers found → skip embedded to avoid the partial-overlay + // mixing bug (ROCm-issue-2445). + std::string FoundPath; + if (detectSystemCxxHeadersOnDisk(Argv, &FoundPath)) + return Decide(true, Twine("system C++ headers found at ") + FoundPath); + return Decide(false, "no system C++ headers found, falling back to embedded"); +} + amd_comgr_status_t AMDGPUCompiler::processFile(DataObject *Input, const char *InputFilePath, const char *OutputFilePath) { @@ -1049,14 +1170,18 @@ amd_comgr_status_t AMDGPUCompiler::processFile(DataObject *Input, Argv.push_back("-nogpulib"); } - // Auto-inject embedded libc++ headers as a fallback include path. - // Using -idirafter places them AFTER all other include paths, so: - // - System libstdc++ or libc++ headers take priority when available - // - User-provided -I paths take priority - // - Embedded headers only kick in when no other C++ headers are found - // This ensures backward compatibility while providing headers on systems - // without C++ development headers (e.g., driver-only installs). - if (HasEmbeddedHeaders && getLanguage() == AMD_COMGR_LANGUAGE_HIP) { + // Auto-inject embedded libc++ headers as a fallback include path when no + // system C++ headers are available. We deliberately skip injection when the + // host has libstdc++ or libc++ installed, because the embedded set is + // partial (LIBCXX_USER_HEADERS in cmake/LibcxxHeaders.cmake) — leaving the + // host's headers in the search path alongside an `-idirafter` to embedded + // libc++ produces a hybrid include chain that breaks `#include_next` + // resolution from libstdc++ (ROCm-issue-2445). + // + // Using -idirafter places them AFTER all other include paths, so when we + // do inject, user `-I`, system headers, and `-isystem` still take priority. + if (HasEmbeddedHeaders && getLanguage() == AMD_COMGR_LANGUAGE_HIP && + !shouldSkipEmbeddedHeaders(Argv)) { SmallString<256> LibcxxPath(env::getLLVMPath()); sys::path::append(LibcxxPath, "include", "c++", "v1"); Argv.push_back("-idirafter"); @@ -1314,10 +1439,7 @@ amd_comgr_status_t AMDGPUCompiler::outputResource(llvm::StringRef Path, } amd_comgr_status_t AMDGPUCompiler::addDeviceLibraries() { - SmallString<256> ClangBinaryPath(env::getLLVMPath()); - sys::path::append(ClangBinaryPath, "bin", "clang"); - - std::string ClangResourceDir = GetResourcesPath(ClangBinaryPath); + std::string ClangResourceDir = GetResourcesPath(env::getClangBinaryPath()); NoGpuLib = false; @@ -2417,9 +2539,7 @@ AMDGPUCompiler::AMDGPUCompiler(DataAction *ActionInfo, DataSet *InSet, OverlayFS->pushOverlay(InMemoryFS); } - SmallString<256> ClangBinaryPath(env::getLLVMPath()); - sys::path::append(ClangBinaryPath, "bin", "clang"); - std::string ResourceDir = GetResourcesPath(ClangBinaryPath); + std::string ResourceDir = GetResourcesPath(env::getClangBinaryPath()); // libc++ headers → /include/c++/v1/ SmallString<256> LibcxxBase(env::getLLVMPath()); diff --git a/amd/comgr/src/comgr-compiler.h b/amd/comgr/src/comgr-compiler.h index 0a1cee6d2db3d..8d024eb5d85ed 100644 --- a/amd/comgr/src/comgr-compiler.h +++ b/amd/comgr/src/comgr-compiler.h @@ -40,6 +40,9 @@ class AMDGPUCompiler { bool UseVFS = false; /// Whether embedded libc++ headers were loaded into the VFS. bool HasEmbeddedHeaders = false; + /// Cached result of `shouldSkipEmbeddedHeaders`, computed once per compiler + /// instance from user args + filesystem probe + env override. + std::optional SkipEmbeddedHeadersCache; llvm::IntrusiveRefCntPtr OverlayFS; llvm::IntrusiveRefCntPtr InMemoryFS; @@ -64,6 +67,12 @@ class AMDGPUCompiler { amd_comgr_status_t executeInProcessDriver(llvm::ArrayRef Args); + /// Decide whether to bypass embedded libc++ headers (skip `-idirafter` + /// injection) for this compilation. Returns true when system C++ headers + /// are available, when the user passed `-nostdinc++`, or when overridden + /// via AMD_COMGR_USE_EMBEDDED_LIBCXX=disable. Cached after first call. + bool shouldSkipEmbeddedHeaders(llvm::ArrayRef Argv); + amd_comgr_status_t translateSpirvToBitcodeImpl(DataSet *SpirvInSet, DataSet *BcOutSet); diff --git a/amd/comgr/src/comgr-env.cpp b/amd/comgr/src/comgr-env.cpp index bd566faa52636..6a629a18d34eb 100644 --- a/amd/comgr/src/comgr-env.cpp +++ b/amd/comgr/src/comgr-env.cpp @@ -14,8 +14,14 @@ #include "comgr-env.h" #include "llvm/ADT/Twine.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/Path.h" #include "llvm/Support/VirtualFileSystem.h" +#ifndef _WIN32 +#include +#endif + using namespace llvm; namespace COMGR { @@ -69,6 +75,59 @@ llvm::StringRef getLLVMPath() { return EnvLLVMPath; } +// Probe whether path P names a clang binary whose derived resource directory +// (parent_path(parent_path(P)) / "lib" / "clang" / ) exists on disk. +// We don't require the binary itself to be executable — clang's Driver only +// uses the path to derive the resource dir. +static bool probeClangResourceDir(StringRef P) { + SmallString<256> ResourceDir(sys::path::parent_path(sys::path::parent_path(P))); + // CLANG_INSTALL_LIBDIR_BASENAME is "lib" by default; we don't link clang, so + // hardcode the common case here. The version directory is what matters for + // existence — its parent will exist on any layout we care about. + sys::path::append(ResourceDir, "lib", "clang"); + return sys::fs::is_directory(ResourceDir); +} + +std::string getClangBinaryPath() { + // Cache the resolved path: this is called from multiple sites per + // compilation and the resolution involves filesystem probes. + static const std::string Cached = []() -> std::string { + // 1. Honor LLVM_PATH explicitly when set. + if (!getLLVMPath().empty()) + return (Twine(getLLVMPath()) + "/bin/clang").str(); + +#ifndef _WIN32 + // 2. Locate libamd_comgr via dladdr and probe sibling layouts. Comgr is + // a shared library; argv[0] of the host process is unrelated to where + // clang lives. dladdr on any symbol in this translation unit yields the + // path of the loaded .so. + Dl_info Info; + if (dladdr(reinterpret_cast(&getClangBinaryPath), &Info) && + Info.dli_fname) { + StringRef SoDir = sys::path::parent_path(Info.dli_fname); + // ROCm packaging: /lib/libamd_comgr.so + /llvm/bin/clang + SmallString<256> RocmLayout(sys::path::parent_path(SoDir)); + sys::path::append(RocmLayout, "llvm", "bin", "clang"); + if (probeClangResourceDir(RocmLayout)) + return std::string(RocmLayout); + + // Standard install: /lib/libamd_comgr.so + /bin/clang + SmallString<256> StandardLayout(sys::path::parent_path(SoDir)); + sys::path::append(StandardLayout, "bin", "clang"); + if (probeClangResourceDir(StandardLayout)) + return std::string(StandardLayout); + } +#endif + + // 3. Fallback: synthesize an absolute "/bin/clang". Resource-dir lookup + // will resolve to "/lib/clang/", which won't exist on disk but + // matches what comgr's VFS embeds — keeping Driver and VFS in sync even + // when the install layout can't be located. + return std::string("/bin/clang"); + }(); + return Cached; +} + StringRef getCachePolicy() { static const char *EnvCachePolicy = std::getenv("AMD_COMGR_CACHE_POLICY"); return EnvCachePolicy; @@ -103,5 +162,17 @@ StringRef getDriverOptionsAppend() { return Options ? Options : ""; } +EmbeddedLibcxxMode getEmbeddedLibcxxMode() { + static const char *V = std::getenv("AMD_COMGR_USE_EMBEDDED_LIBCXX"); + if (!V) + return EmbeddedLibcxxMode::Auto; + StringRef S(V); + if (S.equals_insensitive("force") || S == "1") + return EmbeddedLibcxxMode::Force; + if (S.equals_insensitive("disable") || S == "0") + return EmbeddedLibcxxMode::Disable; + return EmbeddedLibcxxMode::Auto; +} + } // namespace env } // namespace COMGR diff --git a/amd/comgr/src/comgr-env.h b/amd/comgr/src/comgr-env.h index e198a1c67c78a..61f33e917e3c4 100644 --- a/amd/comgr/src/comgr-env.h +++ b/amd/comgr/src/comgr-env.h @@ -33,6 +33,15 @@ bool needTimeStatistics(); /// otherwise return the default LLVM path. llvm::StringRef getLLVMPath(); +/// Return the clang binary path "/bin/clang", constructed via +/// Twine concatenation. This matches the path passed to clang's Driver +/// constructor so that clang::GetResourcesPath() yields a resource-dir path +/// matching what comgr embeds into the VFS, regardless of whether LLVM_PATH +/// is set. Using sys::path::append on an empty base produces a relative +/// "bin/clang" instead of "/bin/clang", which would cause the resource dir +/// to drift between Driver and embed code. +std::string getClangBinaryPath(); + /// If environment variable AMD_COMGR_CACHE_POLICY is set, return the /// environment variable, otherwise return empty llvm::StringRef getCachePolicy(); @@ -46,6 +55,15 @@ llvm::StringRef getCacheDirectory(); /// space-separated options to append to clang driver invocations. llvm::StringRef getDriverOptionsAppend(); +/// Override for embedded libc++ header injection. +/// Auto — detect system C++ headers and skip embedded if found (default). +/// Force — always inject embedded headers, ignore detection. +/// Disable — never inject embedded headers, regardless of detection. +enum class EmbeddedLibcxxMode { Auto, Force, Disable }; + +/// Read AMD_COMGR_USE_EMBEDDED_LIBCXX. Defaults to Auto. +EmbeddedLibcxxMode getEmbeddedLibcxxMode(); + } // namespace env } // namespace COMGR diff --git a/amd/comgr/test/CMakeLists.txt b/amd/comgr/test/CMakeLists.txt index 8ba3d5e7aa898..98abe5fce4740 100644 --- a/amd/comgr/test/CMakeLists.txt +++ b/amd/comgr/test/CMakeLists.txt @@ -247,5 +247,8 @@ add_comgr_test(name_expression_map_test c) add_comgr_test(compile_hip_test c) add_comgr_test(compile_hip_to_relocatable c) add_comgr_test(compile_hip_with_libcxx_test c) +add_comgr_test(compile_hip_stdlibcxx_conflict_test c) +add_comgr_test(compile_hip_with_system_libcxx_test c) +add_comgr_test(compile_hip_distroless_test c) add_comgr_test(mangled_names_hip_test c) #add_comgr_test(unbundle_hip_test c) diff --git a/amd/comgr/test/compile_hip_distroless_test.c b/amd/comgr/test/compile_hip_distroless_test.c new file mode 100644 index 0000000000000..7d96a7aba8680 --- /dev/null +++ b/amd/comgr/test/compile_hip_distroless_test.c @@ -0,0 +1,108 @@ +//===- compile_hip_distroless_test.c --------------------------------------===// +// +// Part of Comgr, under the Apache License v2.0 with LLVM Exceptions. See +// amd/comgr/LICENSE.TXT in this repository for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Verify that the embedded libc++ injection path still works when forced via +// `AMD_COMGR_USE_EMBEDDED_LIBCXX=force`. This simulates a distroless / driver- +// only host where no system C++ headers exist — the env override bypasses the +// detection probe and exercises the same code path that runs on a truly bare +// system. +// +// Source uses only headers from the embedded subset (LIBCXX_USER_HEADERS in +// cmake/LibcxxHeaders.cmake): type_traits, limits, tuple, cstdint, cstddef, +// initializer_list, concepts. Headers outside that subset would fail because +// the embedded set is partial. +// +//===----------------------------------------------------------------------===// + +#include "amd_comgr.h" +#include "common.h" +#include +#include +#include + +#if defined(_WIN32) || defined(_WIN64) +// MSVC: setenv is POSIX. _putenv_s is the closest equivalent and ignores +// the overwrite flag (always overwrites), which matches our usage here. +#define setenv(name, value, overwrite) _putenv_s((name), (value)) +#endif + +const char *HipSource = + "#define __global__ __attribute__((global))\n" + "#define __device__ __attribute__((device))\n" + "\n" + "#include \n" + "#include \n" + "#include \n" + "#include \n" + "\n" + "static_assert(std::is_integral::value, \"int is integral\");\n" + "static_assert(std::numeric_limits::digits == 31,\n" + " \"int32 digits\");\n" + "static_assert(std::tuple_size>::value == 2,\n" + " \"tuple size\");\n" + "\n" + "extern \"C\" __global__ void test_kernel(int *out) {\n" + " std::tuple t{42, 3.14f};\n" + " out[0] = std::get<0>(t);\n" + "}\n"; + +int main(int Argc, char *Argv[]) { + // Force the embedded path even when system C++ headers are present, to + // mirror what comgr does on a distroless host without invasive sysroot setup. + setenv("AMD_COMGR_USE_EMBEDDED_LIBCXX", "force", 1); + + amd_comgr_data_t DataSource; + amd_comgr_data_set_t DataSetIn, DataSetBc; + amd_comgr_action_info_t ActionInfo; + amd_comgr_status_t Status; + + const char *CompileOptions[] = {"-std=c++17", "-nogpuinc"}; + size_t CompileOptionsCount = + sizeof(CompileOptions) / sizeof(CompileOptions[0]); + + Status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_SOURCE, &DataSource); + checkError(Status, "amd_comgr_create_data"); + Status = amd_comgr_set_data(DataSource, strlen(HipSource), HipSource); + checkError(Status, "amd_comgr_set_data"); + Status = amd_comgr_set_data_name(DataSource, "test_distroless.hip"); + checkError(Status, "amd_comgr_set_data_name"); + + Status = amd_comgr_create_data_set(&DataSetIn); + checkError(Status, "amd_comgr_create_data_set"); + Status = amd_comgr_data_set_add(DataSetIn, DataSource); + checkError(Status, "amd_comgr_data_set_add"); + + Status = amd_comgr_create_action_info(&ActionInfo); + checkError(Status, "amd_comgr_create_action_info"); + Status = + amd_comgr_action_info_set_language(ActionInfo, AMD_COMGR_LANGUAGE_HIP); + checkError(Status, "amd_comgr_action_info_set_language"); + Status = amd_comgr_action_info_set_isa_name(ActionInfo, + "amdgcn-amd-amdhsa--gfx906"); + checkError(Status, "amd_comgr_action_info_set_isa_name"); + Status = amd_comgr_action_info_set_option_list(ActionInfo, CompileOptions, + CompileOptionsCount); + checkError(Status, "amd_comgr_action_info_set_option_list"); + + Status = amd_comgr_create_data_set(&DataSetBc); + checkError(Status, "amd_comgr_create_data_set"); + + Status = amd_comgr_do_action( + AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC, ActionInfo, + DataSetIn, DataSetBc); + checkError(Status, "amd_comgr_do_action (compile to BC)"); + + printf("compile_hip_distroless_test PASSED\n"); + + amd_comgr_destroy_action_info(ActionInfo); + amd_comgr_release_data(DataSource); + amd_comgr_destroy_data_set(DataSetIn); + amd_comgr_destroy_data_set(DataSetBc); + + return 0; +} diff --git a/amd/comgr/test/compile_hip_stdlibcxx_conflict_test.c b/amd/comgr/test/compile_hip_stdlibcxx_conflict_test.c new file mode 100644 index 0000000000000..cf5b561ef3fcb --- /dev/null +++ b/amd/comgr/test/compile_hip_stdlibcxx_conflict_test.c @@ -0,0 +1,252 @@ +//===- compile_hip_stdlibcxx_conflict_test.c ------------------------------===// +// +// Part of Comgr, under the Apache License v2.0 with LLVM Exceptions. See +// amd/comgr/LICENSE.TXT in this repository for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Regression test for: embedded libc++ headers conflict with system libstdc++ +// on RHEL/manylinux environments (ROCm/llvm-project#2445). +// +// On systems where gcc-toolset libstdc++ is installed at clang's default search +// path, comgr's -idirafter injection of embedded libc++ causes a mixed-header +// chain: +// system libstdc++ / -> libc++ stddef.h -> #include_next +// fails under -nogpuinc (no next stddef.h on device path). +// +// This test compiles HIP source that includes headers known to trigger the +// conflict (, , ) without passing -nostdinc++. +// It must succeed: either the system does not have conflicting libstdc++ +// headers, or the comgr fix (auto-injecting -nostdinc++ when -nogpuinc is +// active and embedded headers are present) prevents the chain. +// +//===----------------------------------------------------------------------===// + +#include "amd_comgr.h" +#include "common.h" +#include +#include +#include + +#if defined(_WIN32) || defined(_WIN64) +// The libstdc++/libc++ header conflict only manifests on Linux distributions +// shipping libstdc++ at clang's default search paths (RHEL/manylinux/Ubuntu). +// Windows has no equivalent toolchain layout, so the entire scenario is N/A. +int main(int Argc, char *Argv[]) { + printf("compile_hip_stdlibcxx_conflict_test SKIPPED (not applicable on " + "Windows)\n"); + return 0; +} +#else + +#include +#include + +static int fileExists(const char *Path) { + struct stat St; + return stat(Path, &St) == 0; +} + +// Probe whether a clang resource directory containing builtin headers +// (specifically stdarg.h) is reachable on disk at one of the layouts comgr +// auto-detects. Without one, libstdc++'s #include +// chain has nothing to resolve to (the embedded VFS subset only carries +// stddef.h derivatives). In that environment the conflict test cannot +// distinguish "libcxx-skip is broken" from "no clang builtins on disk". +static int hasClangBuiltinHeadersOnDisk(void) { + // Honor LLVM_PATH first if set. + const char *LLVMPath = getenv("LLVM_PATH"); + const char *Roots[] = {LLVMPath ? LLVMPath : "", + "/opt/rocm/llvm", + "/usr", + NULL}; + for (int I = 0; Roots[I]; ++I) { + if (Roots[I][0] == '\0') + continue; + char ClangDir[512]; + snprintf(ClangDir, sizeof(ClangDir), "%s/lib/clang", Roots[I]); + DIR *D = opendir(ClangDir); + if (!D) + continue; + struct dirent *E; + while ((E = readdir(D)) != NULL) { + if (E->d_name[0] == '.') + continue; + char Probe[1024]; + snprintf(Probe, sizeof(Probe), "%s/%s/include/stdarg.h", ClangDir, + E->d_name); + if (fileExists(Probe)) { + closedir(D); + return 1; + } + } + closedir(D); + } + return 0; +} + +// Mirror of detectSystemCxxHeadersOnDisk in comgr-compiler.cpp: returns true +// when either system libc++ or libstdc++ headers are installed at standard +// paths. Used to skip this test on truly bare hosts (no system C++ headers +// at all), where the embedded subset cannot resolve /. +static int hasSystemCxxHeaders(void) { + if (fileExists("/usr/include/c++/v1/__config_site") || + fileExists("/usr/local/include/c++/v1/__config_site")) + return 1; + DIR *D = opendir("/usr/include/c++"); + if (!D) + return 0; + struct dirent *E; + int Found = 0; + while ((E = readdir(D)) != NULL) { + if (E->d_name[0] == '.') + continue; + char Probe[512]; + snprintf(Probe, sizeof(Probe), "/usr/include/c++/%s/cstddef", E->d_name); + if (fileExists(Probe)) { + Found = 1; + break; + } + } + closedir(D); + return Found; +} + +// HIP source using headers that trigger the libstdc++/libc++ conflict. +// , , all transitively pull in +// via libstdc++ on RHEL/manylinux/Ubuntu, which then collides with libc++'s +// VFS-mapped stddef.h doing #include_next with no successor under -nogpuinc. +// +// We deliberately do NOT throw on device — devices can't throw. The point is +// that #including these host C++ headers must not break parsing. +const char *HipSource = + "#define __global__ __attribute__((global))\n" + "#define __device__ __attribute__((device))\n" + "\n" + "#include \n" + "#include \n" + "#include \n" + "\n" + "static_assert(std::tuple_size>::value == 3,\n" + " \"array size\");\n" + "\n" + "extern \"C\" __global__ void test_kernel(int *out) {\n" + " std::array a = {1, 2, 3};\n" + " out[0] = a[0] + a[1] + a[2];\n" + "}\n"; + +// Print log data from a data set for diagnostics on failure. +static void printLogs(amd_comgr_data_set_t DataSet) { + size_t Count; + amd_comgr_status_t Status = + amd_comgr_action_data_count(DataSet, AMD_COMGR_DATA_KIND_LOG, &Count); + if (Status != AMD_COMGR_STATUS_SUCCESS) + return; + for (size_t i = 0; i < Count; i++) { + amd_comgr_data_t Data; + Status = amd_comgr_action_data_get_data(DataSet, AMD_COMGR_DATA_KIND_LOG, i, + &Data); + if (Status != AMD_COMGR_STATUS_SUCCESS) + continue; + size_t Size; + Status = amd_comgr_get_data(Data, &Size, NULL); + if (Status != AMD_COMGR_STATUS_SUCCESS) { + amd_comgr_release_data(Data); + continue; + } + char *Bytes = (char *)malloc(Size + 1); + if (!Bytes) { + amd_comgr_release_data(Data); + continue; + } + Status = amd_comgr_get_data(Data, &Size, Bytes); + if (Status == AMD_COMGR_STATUS_SUCCESS) { + Bytes[Size] = '\0'; + fprintf(stderr, "comgr log:\n%s\n", Bytes); + } + free(Bytes); + amd_comgr_release_data(Data); + } +} + +int main(int Argc, char *Argv[]) { + // // are not in the embedded libc++ subset, + // so on a host with neither libstdc++ nor libc++ this test cannot + // distinguish "fix is broken" from "no host C++ stdlib at all". Skip. + if (!hasSystemCxxHeaders()) { + printf("compile_hip_stdlibcxx_conflict_test SKIPPED " + "(no system C++ headers found)\n"); + return 0; + } + + // If no clang resource dir is reachable on disk, libstdc++'s + // -> chain cannot resolve regardless of comgr's + // libcxx-skip behavior. Skip rather than report a misleading failure. + if (!hasClangBuiltinHeadersOnDisk()) { + printf("compile_hip_stdlibcxx_conflict_test SKIPPED " + "(no clang builtin headers on disk; set LLVM_PATH or install " + "clang at /opt/rocm/llvm or /usr)\n"); + return 0; + } + + amd_comgr_data_t DataSource; + amd_comgr_data_set_t DataSetIn, DataSetBc; + amd_comgr_action_info_t ActionInfo; + amd_comgr_status_t Status; + + // No -nostdinc++ here: that is the point of this test. Comgr must handle + // the libstdc++/libc++ conflict internally (e.g. by auto-injecting + // -nostdinc++ when -nogpuinc is active and embedded headers are present). + const char *CompileOptions[] = {"-std=c++17", "-nogpuinc"}; + size_t CompileOptionsCount = + sizeof(CompileOptions) / sizeof(CompileOptions[0]); + + Status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_SOURCE, &DataSource); + checkError(Status, "amd_comgr_create_data"); + Status = amd_comgr_set_data(DataSource, strlen(HipSource), HipSource); + checkError(Status, "amd_comgr_set_data"); + Status = amd_comgr_set_data_name(DataSource, "test_conflict.hip"); + checkError(Status, "amd_comgr_set_data_name"); + + Status = amd_comgr_create_data_set(&DataSetIn); + checkError(Status, "amd_comgr_create_data_set"); + Status = amd_comgr_data_set_add(DataSetIn, DataSource); + checkError(Status, "amd_comgr_data_set_add"); + + Status = amd_comgr_create_action_info(&ActionInfo); + checkError(Status, "amd_comgr_create_action_info"); + Status = + amd_comgr_action_info_set_language(ActionInfo, AMD_COMGR_LANGUAGE_HIP); + checkError(Status, "amd_comgr_action_info_set_language"); + Status = amd_comgr_action_info_set_isa_name(ActionInfo, + "amdgcn-amd-amdhsa--gfx906"); + checkError(Status, "amd_comgr_action_info_set_isa_name"); + Status = amd_comgr_action_info_set_option_list(ActionInfo, CompileOptions, + CompileOptionsCount); + checkError(Status, "amd_comgr_action_info_set_option_list"); + + Status = amd_comgr_create_data_set(&DataSetBc); + checkError(Status, "amd_comgr_create_data_set"); + + Status = amd_comgr_do_action( + AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC, ActionInfo, + DataSetIn, DataSetBc); + if (Status != AMD_COMGR_STATUS_SUCCESS) { + printLogs(DataSetBc); + fail("amd_comgr_do_action (compile to BC) -- " + "likely libstdc++/libc++ header conflict (ROCm#2445): " + "comgr must auto-inject -nostdinc++ when -nogpuinc is active"); + } + + printf("compile_hip_stdlibcxx_conflict_test PASSED\n"); + + amd_comgr_destroy_action_info(ActionInfo); + amd_comgr_release_data(DataSource); + amd_comgr_destroy_data_set(DataSetIn); + amd_comgr_destroy_data_set(DataSetBc); + + return 0; +} + +#endif // !_WIN32 diff --git a/amd/comgr/test/compile_hip_with_system_libcxx_test.c b/amd/comgr/test/compile_hip_with_system_libcxx_test.c new file mode 100644 index 0000000000000..9d5800196b1bd --- /dev/null +++ b/amd/comgr/test/compile_hip_with_system_libcxx_test.c @@ -0,0 +1,157 @@ +//===- compile_hip_with_system_libcxx_test.c ------------------------------===// +// +// Part of Comgr, under the Apache License v2.0 with LLVM Exceptions. See +// amd/comgr/LICENSE.TXT in this repository for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Verify that when system libc++ is installed at the standard path +// (`/usr/include/c++/v1/__config_site`), comgr's auto-detection takes the +// libc++ branch of `detectSystemCxxHeadersOnDisk` and refrains from +// `-idirafter` injection. Companion to `compile_hip_stdlibcxx_conflict_test`, +// which covers the libstdc++ branch and the actual end-to-end compile. +// +// This test asserts the detection behavior (verbose log) rather than +// compile success: clang's `-stdlib=libc++` resolution does not redirect to +// `/usr/include/c++/v1`, so a clean compile would require a more invasive +// test setup. The detection branch firing is what matters for the fix. +// +// Skipped on hosts without system libc++ headers (e.g. minimal CI images). +// +//===----------------------------------------------------------------------===// + +#include "amd_comgr.h" +#include "common.h" +#include +#include +#include + +#if defined(_WIN32) || defined(_WIN64) +// The probed paths (/usr/include/c++/v1/__config_site) are Linux-only; +// the Windows toolchain layout has no analogous system libc++ to detect. +int main(int Argc, char *Argv[]) { + printf("compile_hip_with_system_libcxx_test SKIPPED (not applicable on " + "Windows)\n"); + return 0; +} +#else + +#include + +static int fileExists(const char *Path) { + struct stat St; + return stat(Path, &St) == 0; +} + +// Search every LOG entry in DataSet for Needle; fail with id if absent. +static void requireLogContains(const char *Id, amd_comgr_data_set_t DataSet, + const char *Needle) { + size_t Count; + amd_comgr_status_t Status = + amd_comgr_action_data_count(DataSet, AMD_COMGR_DATA_KIND_LOG, &Count); + checkError(Status, "amd_comgr_action_data_count"); + + int Found = 0; + for (size_t I = 0; I < Count && !Found; ++I) { + amd_comgr_data_t Data; + Status = amd_comgr_action_data_get_data(DataSet, AMD_COMGR_DATA_KIND_LOG, I, + &Data); + checkError(Status, "amd_comgr_action_data_get_data"); + + size_t Size; + Status = amd_comgr_get_data(Data, &Size, NULL); + checkError(Status, "amd_comgr_get_data"); + + char *Bytes = (char *)malloc(Size + 1); + if (!Bytes) + fail("malloc"); + Status = amd_comgr_get_data(Data, &Size, Bytes); + checkError(Status, "amd_comgr_get_data"); + Bytes[Size] = '\0'; + + if (strstr(Bytes, Needle)) + Found = 1; + + free(Bytes); + amd_comgr_release_data(Data); + } + + if (!Found) + fail("%s: expected log substring \"%s\" not found", Id, Needle); +} + +const char *HipSource = + "#define __global__ __attribute__((global))\n" + "#define __device__ __attribute__((device))\n" + "\n" + "extern \"C\" __global__ void test_kernel(int *out) { out[0] = 1; }\n"; + +int main(int Argc, char *Argv[]) { + if (!fileExists("/usr/include/c++/v1/__config_site") && + !fileExists("/usr/local/include/c++/v1/__config_site")) { + printf("compile_hip_with_system_libcxx_test SKIPPED " + "(system libc++ not installed)\n"); + return 0; + } + + // Need verbose logs for the "Embedded libc++ headers: …" line. + setenv("AMD_COMGR_EMIT_VERBOSE_LOGS", "1", 1); + + amd_comgr_data_t DataSource; + amd_comgr_data_set_t DataSetIn, DataSetBc; + amd_comgr_action_info_t ActionInfo; + amd_comgr_status_t Status; + + const char *CompileOptions[] = {"-std=c++17", "-nogpuinc"}; + size_t CompileOptionsCount = + sizeof(CompileOptions) / sizeof(CompileOptions[0]); + + Status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_SOURCE, &DataSource); + checkError(Status, "amd_comgr_create_data"); + Status = amd_comgr_set_data(DataSource, strlen(HipSource), HipSource); + checkError(Status, "amd_comgr_set_data"); + Status = amd_comgr_set_data_name(DataSource, "test_system_libcxx.hip"); + checkError(Status, "amd_comgr_set_data_name"); + + Status = amd_comgr_create_data_set(&DataSetIn); + checkError(Status, "amd_comgr_create_data_set"); + Status = amd_comgr_data_set_add(DataSetIn, DataSource); + checkError(Status, "amd_comgr_data_set_add"); + + Status = amd_comgr_create_action_info(&ActionInfo); + checkError(Status, "amd_comgr_create_action_info"); + Status = + amd_comgr_action_info_set_language(ActionInfo, AMD_COMGR_LANGUAGE_HIP); + checkError(Status, "amd_comgr_action_info_set_language"); + Status = amd_comgr_action_info_set_isa_name(ActionInfo, + "amdgcn-amd-amdhsa--gfx906"); + checkError(Status, "amd_comgr_action_info_set_isa_name"); + Status = amd_comgr_action_info_set_option_list(ActionInfo, CompileOptions, + CompileOptionsCount); + checkError(Status, "amd_comgr_action_info_set_option_list"); + Status = amd_comgr_action_info_set_logging(ActionInfo, true); + checkError(Status, "amd_comgr_action_info_set_logging"); + + Status = amd_comgr_create_data_set(&DataSetBc); + checkError(Status, "amd_comgr_create_data_set"); + + // Compile outcome is not asserted — we only care that detection ran and + // logged the skip decision. + amd_comgr_do_action(AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC, + ActionInfo, DataSetIn, DataSetBc); + + requireLogContains("compile_hip_with_system_libcxx_test", DataSetBc, + "Embedded libc++ headers: skipped"); + + printf("compile_hip_with_system_libcxx_test PASSED\n"); + + amd_comgr_destroy_action_info(ActionInfo); + amd_comgr_release_data(DataSource); + amd_comgr_destroy_data_set(DataSetIn); + amd_comgr_destroy_data_set(DataSetBc); + + return 0; +} + +#endif // !_WIN32