[clang] [compiler-rt] [compiler-rt] Support `-shared-libsan` for Offload UBSan (PR #226551)
Joseph Huber via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 25 10:46:32 PDT 2026
https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/226551
Summary:
The offload UBSan is intentionally implemented as a separate 'add-on'
library. This makes it more clear where the separation lies and avoids
including it in all existing UBSan users. However, this doesn't work
with `-shared-libsan`'s `*.so` builds because we now have a function
call barrier and the internal utility functions are hidden.
This PR changes the shared version to fold the HSA / Offloading handlers
into the `libclang_rt.ubsan.so` library. The specific fallout is that
~30KiB is added to the shared library and it will now intercept `dlsym`
for other users, but this should be benign. The overhead here is less
concerning than in the static case because it is shared and usually
loaded lazily by the system.
>From 1a7b445e705ba93f5bf00b4eaa1b22ab7f526772 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 25 Sep 2026 12:40:38 -0500
Subject: [PATCH] [compiler-rt] Support `-shared-libsan` for Offload UBSan
Summary:
The offload UBSan is intentionally implemented as a separate 'add-on'
library. This makes it more clear where the separation lies and avoids
including it in all existing UBSan users. However, this doesn't work
with `-shared-libsan`'s `*.so` builds because we now have a function
call barrier and the internal utility functions are hidden.
This PR changes the shared version to fold the HSA / Offloading handlers
into the `libclang_rt.ubsan.so` library. The specific fallout is that
~30KiB is added to the shared library and it will now intercept `dlsym`
for other users, but this should be benign. The overhead here is less
concerning than in the static case because it is shared and usually
loaded lazily by the system.
---
clang/lib/Driver/ToolChains/CommonArgs.cpp | 5 ++--
.../test/Driver/fsanitize-undefined-offload.c | 8 ++++++
compiler-rt/lib/ubsan/CMakeLists.txt | 20 ++++++++++----
compiler-rt/lib/ubsan/offload/CMakeLists.txt | 18 ++++++++++---
.../ubsan_offload_hsa_interceptors.cpp | 26 +++++++++++--------
.../test/ubsan/AMDGPU/host-and-device.hip | 2 ++
6 files changed, 57 insertions(+), 22 deletions(-)
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index 445eb4ccfbfa72..1200784df1f1c9 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -1744,8 +1744,9 @@ collectSanitizerRuntimes(Compilation &C, const ToolChain &TC,
if (SanArgs.needsAsanRt())
HelperStaticRuntimes.push_back("asan_static");
- // Offloading images can live in DSOs, the host interceptors must follow.
- if (NeedsOffloadRt) {
+ // Offloading images can live in DSOs, the host interceptors must follow. The
+ // shared UBSan runtime already contains them.
+ if (NeedsOffloadRt && !SanArgs.needsSharedRt()) {
NonWholeStaticRuntimes.push_back("ubsan_offload");
RequiredSymbols.push_back("__ubsan_offload_init");
}
diff --git a/clang/test/Driver/fsanitize-undefined-offload.c b/clang/test/Driver/fsanitize-undefined-offload.c
index 424ae23f709672..e9f6be10225178 100644
--- a/clang/test/Driver/fsanitize-undefined-offload.c
+++ b/clang/test/Driver/fsanitize-undefined-offload.c
@@ -50,6 +50,14 @@
// CHECK-SHARED-DAG: "-u" "__ubsan_offload_init"
// CHECK-SHARED-DAG: "{{[^"]*}}x86_64-unknown-linux-gnu{{/|\\\\}}libclang_rt.ubsan_offload.a"
+// RUN: %clang -no-canonical-prefixes -### --target=x86_64-unknown-linux-gnu \
+// RUN: -x hip --offload-arch=gfx908 -fsanitize=undefined -shared-libsan \
+// RUN: -nogpuinc -nogpulib --rocm-path=%S/Inputs/rocm \
+// RUN: -resource-dir=%S/Inputs/resource_dir_with_amdgpu_per_target_subdir %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CHECK-SHARED-RT \
+// RUN: --implicit-check-not=ubsan_offload
+// CHECK-SHARED-RT: "{{[^"]*}}x86_64-unknown-linux-gnu{{/|\\\\}}libclang_rt.ubsan_standalone.so"
+
// RUN: %clang -no-canonical-prefixes -### --target=x86_64-unknown-linux-gnu \
// RUN: -x hip --offload-arch=gfx908 -Xarch_gfx908 -fsanitize=undefined \
// RUN: -nogpuinc -nogpulib --rocm-path=%S/Inputs/rocm \
diff --git a/compiler-rt/lib/ubsan/CMakeLists.txt b/compiler-rt/lib/ubsan/CMakeLists.txt
index f49040222ac50b..d03452ea983a90 100644
--- a/compiler-rt/lib/ubsan/CMakeLists.txt
+++ b/compiler-rt/lib/ubsan/CMakeLists.txt
@@ -213,6 +213,12 @@ else()
DEFS ${UBSAN_COMMON_DEFINITIONS})
endif()
+ # Thin host interceptors for offloading. Linked with -u to keep it separate
+ # from static runtimes and embedded directly in the shared runtime.
+ if(OS_NAME MATCHES "Linux" AND NOT ANDROID)
+ add_subdirectory(offload)
+ endif()
+
if(COMPILER_RT_HAS_UBSAN)
add_compiler_rt_object_libraries(RTUbsan_standalone
ARCHS ${UBSAN_SUPPORTED_ARCH}
@@ -260,9 +266,17 @@ else()
CFLAGS ${UBSAN_CFLAGS})
foreach(arch ${UBSAN_SUPPORTED_ARCH})
+ set(UBSAN_OFFLOAD_OBJECT_LIBS)
+ set(UBSAN_OFFLOAD_VERSION_LIBS)
+ if(TARGET RTUbsan_offload.${arch} AND TARGET RTSanitizerOffload.${arch})
+ set(UBSAN_OFFLOAD_OBJECT_LIBS RTUbsan_offload RTSanitizerOffload)
+ set(UBSAN_OFFLOAD_VERSION_LIBS clang_rt.ubsan_offload-${arch})
+ endif()
+
add_sanitizer_rt_version_list(clang_rt.ubsan_standalone-dynamic-${arch}
LIBS clang_rt.ubsan_standalone-${arch}
clang_rt.ubsan_standalone_cxx-${arch}
+ ${UBSAN_OFFLOAD_VERSION_LIBS}
EXTRA ubsan.syms.extra)
set(VERSION_SCRIPT_FLAG
-Wl,--version-script,${CMAKE_CURRENT_BINARY_DIR}/clang_rt.ubsan_standalone-dynamic-${arch}.vers)
@@ -288,6 +302,7 @@ else()
RTUbsan_cxx
RTUbsan_standalone
RTInterception
+ ${UBSAN_OFFLOAD_OBJECT_LIBS}
RTUbsan_dynamic_version_script_dummy
CFLAGS ${UBSAN_CFLAGS}
LINK_FLAGS ${UBSAN_LINK_FLAGS} ${VERSION_SCRIPT_FLAG}
@@ -311,9 +326,4 @@ else()
EXTRA ubsan.syms.extra)
endif()
endif()
-
- # Thin host interceptors for offloading. Linked with -u to keep it separate.
- if(OS_NAME MATCHES "Linux" AND NOT ANDROID)
- add_subdirectory(offload)
- endif()
endif()
diff --git a/compiler-rt/lib/ubsan/offload/CMakeLists.txt b/compiler-rt/lib/ubsan/offload/CMakeLists.txt
index b6046350d907e7..6f03512ef2ba9b 100644
--- a/compiler-rt/lib/ubsan/offload/CMakeLists.txt
+++ b/compiler-rt/lib/ubsan/offload/CMakeLists.txt
@@ -15,12 +15,22 @@ set(UBSAN_OFFLOAD_HEADERS
../ubsan_handlers_internal.h
)
-add_compiler_rt_runtime(clang_rt.ubsan_offload
- STATIC
+add_compiler_rt_object_libraries(RTUbsan_offload
ARCHS ${UBSAN_SUPPORTED_ARCH}
SOURCES ${UBSAN_OFFLOAD_SOURCES}
ADDITIONAL_HEADERS ${UBSAN_OFFLOAD_HEADERS}
- OBJECT_LIBS RTSanitizerOffload
+ CFLAGS ${UBSAN_CFLAGS})
+foreach(arch ${UBSAN_SUPPORTED_ARCH})
+ target_link_libraries(RTUbsan_offload.${arch} PRIVATE
+ llvm-libc-common-utilities)
+endforeach()
+
+# The shared runtime embeds these objects directly, the static archive is only
+# for static runtimes and is pulled in with -u.
+add_compiler_rt_runtime(clang_rt.ubsan_offload
+ STATIC
+ ARCHS ${UBSAN_SUPPORTED_ARCH}
+ OBJECT_LIBS RTUbsan_offload
+ RTSanitizerOffload
CFLAGS ${UBSAN_CFLAGS}
- LINK_LIBS llvm-libc-common-utilities
PARENT_TARGET ubsan)
diff --git a/compiler-rt/lib/ubsan/offload/ubsan_offload_hsa_interceptors.cpp b/compiler-rt/lib/ubsan/offload/ubsan_offload_hsa_interceptors.cpp
index 5f251c2fc2a3ab..81dca67fe99e5e 100644
--- a/compiler-rt/lib/ubsan/offload/ubsan_offload_hsa_interceptors.cpp
+++ b/compiler-rt/lib/ubsan/offload/ubsan_offload_hsa_interceptors.cpp
@@ -67,8 +67,15 @@ void Initialize() {
if (UNLIKELY(!Offload::Get().Ready())) \
return REAL(name)(__VA_ARGS__);
-// PPC cannot transparently tail-call an indirect dlsym target for RTLD_NEXT.
-#if !SANITIZER_PPC
+// The dlsym interceptor is only transparent if it can tail-call, which PPC and
+// compilers without musttail cannot guarantee.
+#if defined(__has_cpp_attribute)
+#if __has_cpp_attribute(clang::musttail) && !SANITIZER_PPC
+#define UBSAN_INTERCEPT_DLSYM 1
+#endif
+#endif
+
+#if UBSAN_INTERCEPT_DLSYM
#define UBSAN_HSA_WRAPS(X) \
X(hsa_init) \
X(hsa_shut_down) \
@@ -96,20 +103,17 @@ static void BindRealDlsym();
// OpenMP and sometimes HIP access HSA through 'dlsym' so we need to intercept
// it here if we want to reliably override its definitions.
INTERCEPTOR(void *, dlsym, void *Handle, const char *Name) {
- Initialize();
BindRealDlsym();
- // This interceptor interferes with the order of 'RTLD_NEXT'. Force a tail
- // call to bypass this process in the stack.
- if (Handle == RTLD_NEXT) [[clang::musttail]]
+ // glibc resolves 'RTLD_NEXT' and 'RTLD_DEFAULT' relative to the caller's
+ // scope and attributes the new dependency to it. Tail call to preserve it.
+ void *Wrapper = Name ? WrapperFor(Name) : nullptr;
+ if (Handle == RTLD_NEXT || !Wrapper) [[clang::musttail]]
return REAL(dlsym)(Handle, Name);
+ Initialize();
void *Sym = REAL(dlsym)(Handle, Name);
- if (!Sym || !Name)
- return Sym;
-
- void *Wrapper = WrapperFor(Name);
- if (!Wrapper || !FromHsa(Sym))
+ if (!Sym || !FromHsa(Sym))
return Sym;
return Wrapper;
}
diff --git a/compiler-rt/test/ubsan/AMDGPU/host-and-device.hip b/compiler-rt/test/ubsan/AMDGPU/host-and-device.hip
index eb896461bd2c8a..bebadad755ffd1 100644
--- a/compiler-rt/test/ubsan/AMDGPU/host-and-device.hip
+++ b/compiler-rt/test/ubsan/AMDGPU/host-and-device.hip
@@ -1,6 +1,8 @@
// REQUIRES: ubsan-hip
// RUN: %clang_hip -fsanitize=undefined %s -o %t %hip_libs
// RUN: %run %t 2>&1 | FileCheck %s
+// RUN: %clang_hip -fsanitize=undefined -shared-libsan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
__global__ void gpu_shift(int *Out) {
int A = 1;
More information about the llvm-commits
mailing list