From eab515b6d9519dfab6866aab0185b51c183a0390 Mon Sep 17 00:00:00 2001 From: Victor Mustya Date: Tue, 5 May 2026 17:46:18 -0700 Subject: [PATCH 1/4] [compiler-rt] Support target-specific handlers for ubsan_minimal runtime Some targets, e.g. GPU ones, may require custom implementations of UBSan message printing and aborting functions, because the generic implementation may be not possible to support on such targers. This patch adds support for target-specific handlers in the minimal UBSan runtime, and outlines the existing generic implementation as a fallback for targets that do not require custom handlers. The patch also adjusts the CMake configuration for simpler integration of GPU targets support in the minimal UBSan runtime. --- compiler-rt/CMakeLists.txt | 7 +- .../cmake/Modules/AllSupportedArchDefs.cmake | 4 +- compiler-rt/cmake/config-ix.cmake | 3 + .../lib/sanitizer_common/sanitizer_platform.h | 6 +- compiler-rt/lib/ubsan_minimal/CMakeLists.txt | 34 +++++--- .../lib/ubsan_minimal/targets/generic.cpp | 76 ++++++++++++++++ .../lib/ubsan_minimal/ubsan_minimal_common.h | 32 +++++++ .../ubsan_minimal/ubsan_minimal_handlers.cpp | 87 ++----------------- 8 files changed, 153 insertions(+), 96 deletions(-) create mode 100644 compiler-rt/lib/ubsan_minimal/targets/generic.cpp create mode 100644 compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt index 6e8045631b218..360095f0b88e5 100644 --- a/compiler-rt/CMakeLists.txt +++ b/compiler-rt/CMakeLists.txt @@ -370,7 +370,7 @@ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "hexagon") list(APPEND SANITIZER_COMMON_CFLAGS -fno-emulated-tls) string(APPEND COMPILER_RT_TEST_COMPILER_CFLAGS " -fno-emulated-tls") endif() -if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx") +if(COMPILER_RT_GPU_BUILD) append_list_if(UBSAN_COMMON_SUPPORTED_ARCH -ffreestanding SANITIZER_COMMON_CFLAGS) append_list_if(COMPILER_RT_HAS_NOGPULIB_FLAG -nogpulib SANITIZER_COMMON_CFLAGS) append_list_if(COMPILER_RT_HAS_FLTO_FLAG -flto SANITIZER_COMMON_CFLAGS) @@ -378,6 +378,8 @@ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx") append_list_if(COMPILER_RT_HAS_CODE_OBJECT_VERSION_FLAG "SHELL:-Xclang -mcode-object-version=none" SANITIZER_COMMON_CFLAGS) endif() +else() + append_list_if(COMPILER_RT_HAS_FNO_LTO_FLAG -fno-lto SANITIZER_COMMON_CFLAGS) endif() if(NOT WIN32) @@ -395,9 +397,6 @@ append_list_if(COMPILER_RT_HAS_FVISIBILITY_HIDDEN_FLAG -fvisibility=hidden SANIT if(NOT COMPILER_RT_HAS_FVISIBILITY_HIDDEN_FLAG) append_list_if(COMPILER_RT_HAS_FVISIBILITY_INLINES_HIDDEN_FLAG -fvisibility-inlines-hidden SANITIZER_COMMON_CFLAGS) endif() -if(NOT "${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx") - append_list_if(COMPILER_RT_HAS_FNO_LTO_FLAG -fno-lto SANITIZER_COMMON_CFLAGS) -endif() # By default do not instrument or use profdata for compiler-rt. if(NOT COMPILER_RT_ENABLE_PGO) diff --git a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake index 761e26ce3e7d8..fa8b8d6679582 100644 --- a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake +++ b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake @@ -105,7 +105,9 @@ endif() set(ALL_TYSAN_SUPPORTED_ARCH ${X86_64} ${ARM64} ${S390X} ${HEXAGON}) set(ALL_UBSAN_SUPPORTED_ARCH ${X86} ${X86_64} ${ARM32} ${ARM64} ${RISCV64} ${MIPS32} ${MIPS64} ${PPC64} ${S390X} ${SPARC} ${SPARCV9} ${HEXAGON} - ${LOONGARCH64} ${AMDGPU} ${NVPTX}) + ${LOONGARCH64}) +set(ALL_UBSAN_MINIMAL_SUPPORTED_ARCH + ${ALL_UBSAN_SUPPORTED_ARCH} ${AMDGPU} ${NVPTX}) if (OS_NAME MATCHES "FreeBSD") set(ALL_SAFESTACK_SUPPORTED_ARCH ${X86} ${X86_64} ${ARM64}) else() diff --git a/compiler-rt/cmake/config-ix.cmake b/compiler-rt/cmake/config-ix.cmake index 9909783d7bb57..8129fb662c72d 100644 --- a/compiler-rt/cmake/config-ix.cmake +++ b/compiler-rt/cmake/config-ix.cmake @@ -727,6 +727,7 @@ else() filter_available_targets(TSAN_SUPPORTED_ARCH ${ALL_TSAN_SUPPORTED_ARCH}) filter_available_targets(TYSAN_SUPPORTED_ARCH ${ALL_TYSAN_SUPPORTED_ARCH}) filter_available_targets(UBSAN_SUPPORTED_ARCH ${ALL_UBSAN_SUPPORTED_ARCH}) + filter_available_targets(UBSAN_MINIMAL_SUPPORTED_ARCH ${ALL_UBSAN_MINIMAL_SUPPORTED_ARCH}) filter_available_targets(SAFESTACK_SUPPORTED_ARCH ${ALL_SAFESTACK_SUPPORTED_ARCH}) filter_available_targets(CFI_SUPPORTED_ARCH ${ALL_CFI_SUPPORTED_ARCH}) @@ -896,6 +897,8 @@ if (UBSAN_SUPPORTED_ARCH AND (OS_NAME MATCHES "Linux|FreeBSD|NetBSD|Android|Darwin|SunOS" OR "${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx")) set(COMPILER_RT_HAS_UBSAN_MINIMAL TRUE) +elseif(COMPILER_RT_GPU_BUILD AND UBSAN_MINIMAL_SUPPORTED_ARCH) + set(COMPILER_RT_HAS_UBSAN_MINIMAL TRUE) else() set(COMPILER_RT_HAS_UBSAN_MINIMAL FALSE) endif() diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h index a1f3cf799e344..daa62de131fe4 100644 --- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h +++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h @@ -12,10 +12,14 @@ #ifndef SANITIZER_PLATFORM_H #define SANITIZER_PLATFORM_H +#if !defined(SANITIZER_GPU) +# define SANITIZER_GPU 0 +#endif // !defined(SANITIZER_GPU) + #if !defined(__linux__) && !defined(__FreeBSD__) && !defined(__NetBSD__) && \ !defined(__APPLE__) && !defined(_WIN32) && !defined(__Fuchsia__) && \ !(defined(__sun__) && defined(__svr4__)) && !defined(__HAIKU__) && \ - !defined(__wasi__) && !defined(__NVPTX__) && !defined(__AMDGPU__) + !defined(__wasi__) && !(SANITIZER_GPU) # error "This operating system is not supported" #endif diff --git a/compiler-rt/lib/ubsan_minimal/CMakeLists.txt b/compiler-rt/lib/ubsan_minimal/CMakeLists.txt index 1f6e50857e9dd..a35187460d2ef 100644 --- a/compiler-rt/lib/ubsan_minimal/CMakeLists.txt +++ b/compiler-rt/lib/ubsan_minimal/CMakeLists.txt @@ -4,6 +4,13 @@ set(UBSAN_MINIMAL_SOURCES ubsan_minimal_handlers.cpp ) +if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/targets/${UBSAN_MINIMAL_SUPPORTED_ARCH}.cpp) + list(APPEND UBSAN_MINIMAL_SOURCES targets/${UBSAN_MINIMAL_SUPPORTED_ARCH}.cpp) +else() + list(APPEND UBSAN_MINIMAL_SOURCES targets/generic.cpp) +endif() + +include_directories(.) include_directories(..) set(UBSAN_CFLAGS @@ -20,7 +27,8 @@ add_compiler_rt_component(ubsan-minimal) # Common parts of minimal UBSan runtime. add_compiler_rt_object_libraries(RTUbsan_minimal OS ${SANITIZER_COMMON_SUPPORTED_OS} - ARCHS ${UBSAN_COMMON_SUPPORTED_ARCH} + ARCHS ${UBSAN_MINIMAL_SUPPORTED_ARCH} + DEFS $,SANITIZER_GPU=1,SANITIZER_GPU=0> SOURCES ${UBSAN_MINIMAL_SOURCES} CFLAGS ${UBSAN_CFLAGS}) @@ -28,23 +36,25 @@ add_compiler_rt_object_libraries(RTUbsan_minimal add_compiler_rt_runtime(clang_rt.ubsan_minimal STATIC OS ${UBSAN_SUPPORTED_OS} - ARCHS ${UBSAN_SUPPORTED_ARCH} + ARCHS ${UBSAN_MINIMAL_SUPPORTED_ARCH} OBJECT_LIBS RTUbsan_minimal CFLAGS ${UBSAN_CFLAGS} PARENT_TARGET ubsan-minimal) -add_compiler_rt_runtime(clang_rt.ubsan_minimal - SHARED - OS ${UBSAN_SUPPORTED_OS} - ARCHS ${UBSAN_SUPPORTED_ARCH} - OBJECT_LIBS RTUbsan_minimal - CFLAGS ${UBSAN_CFLAGS} - LINK_FLAGS ${UBSAN_LINK_FLAGS} - LINK_LIBS ${UBSAN_DYNAMIC_LIBS} - PARENT_TARGET ubsan-minimal) +if(NOT COMPILER_RT_GPU_BUILD) + add_compiler_rt_runtime(clang_rt.ubsan_minimal + SHARED + OS ${UBSAN_SUPPORTED_OS} + ARCHS ${UBSAN_MINIMAL_SUPPORTED_ARCH} + OBJECT_LIBS RTUbsan_minimal + CFLAGS ${UBSAN_CFLAGS} + LINK_FLAGS ${UBSAN_LINK_FLAGS} + LINK_LIBS ${UBSAN_DYNAMIC_LIBS} + PARENT_TARGET ubsan-minimal) +endif() if (SANITIZER_USE_SYMBOLS AND NOT APPLE) - set(ARCHS_FOR_SYMBOLS ${UBSAN_SUPPORTED_ARCH}) + set(ARCHS_FOR_SYMBOLS ${UBSAN_MINIMAL_SUPPORTED_ARCH}) list(REMOVE_ITEM ARCHS_FOR_SYMBOLS i386 i686) add_sanitizer_rt_symbols(clang_rt.ubsan_minimal ARCHS ${ARCHS_FOR_SYMBOLS} diff --git a/compiler-rt/lib/ubsan_minimal/targets/generic.cpp b/compiler-rt/lib/ubsan_minimal/targets/generic.cpp new file mode 100644 index 0000000000000..41c340a357355 --- /dev/null +++ b/compiler-rt/lib/ubsan_minimal/targets/generic.cpp @@ -0,0 +1,76 @@ + +#include "ubsan_minimal_common.h" + +#include +#include + +#if defined(__ANDROID__) +extern "C" __attribute__((weak)) void android_set_abort_message(const char *); +#endif // defined(__ANDROID__) + +#ifdef KERNEL_USE +extern "C" void ubsan_message(const char *msg); +static void message(const char *msg) { ubsan_message(msg); } +#elif defined(SANITIZER_AMDGPU) || defined(SANITIZER_NVPTX) +#include +static void message(const char *msg) { fprintf(stderr, "%s", msg); } +#else +#include +static void message(const char *msg) { (void)write(2, msg, strlen(msg)); } +#endif + +static char *append_str(const char *s, char *buf, const char *end) { + for (const char *p = s; (buf < end) && (*p != '\0'); ++p, ++buf) + *buf = *p; + return buf; +} + +static char *append_hex(uintptr_t d, char *buf, const char *end) { + // Print the address by nibbles. + for (unsigned shift = sizeof(uintptr_t) * 8; shift && buf < end;) { + shift -= 4; + unsigned nibble = (d >> shift) & 0xf; + *(buf++) = nibble < 10 ? nibble + '0' : nibble - 10 + 'a'; + } + return buf; +} + +static void format_msg(const char *kind, uintptr_t caller, char *buf, + const char *end) { + buf = append_str("ubsan: ", buf, end); + buf = append_str(kind, buf, end); + buf = append_str(" by 0x", buf, end); + buf = append_hex(caller, buf, end); + buf = append_str("\n", buf, end); + if (buf == end) + --buf; // Make sure we don't cause a buffer overflow. + *buf = '\0'; +} + +void __ubsan_message(const char *msg) { + message(msg); +} + +void __ubsan_message(const char *kind, uintptr_t caller) { + char buf[128]; + format_msg(kind, caller, buf, buf + sizeof(buf)); + message(buf); +} + +void __ubsan_abort() { + abort(); +} + +void __ubsan_abort_with_message(const char *kind, uintptr_t caller) { + char buf[128]; + format_msg(kind, caller, buf, buf + sizeof(buf)); + +#if defined(__ANDROID__) + if (&android_set_abort_message) + android_set_abort_message(buf); +#else // defined(__ANDROID__) + message(buf); +#endif // defined(__ANDROID__) + + abort(); +} diff --git a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h new file mode 100644 index 0000000000000..0a665df7e50c3 --- /dev/null +++ b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h @@ -0,0 +1,32 @@ +//===-- ubsan_minimal_common.h ----------------------------------*- 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 +// +//===----------------------------------------------------------------------===// +// +// This file is a part of the minimal UB sanitizer runtime. +// +//===----------------------------------------------------------------------===// + +#ifndef UBSAN_MINIMAL_COMMON_H +#define UBSAN_MINIMAL_COMMON_H + +#if defined(__UINTPTR_TYPE__) +typedef __UINTPTR_TYPE__ uintptr_t; + +static_assert(sizeof(uintptr_t) == sizeof(void *), + "uintptr_t must be the same size as void*"); +#else // defined(__UINTPTR_TYPE__) +#include +#endif // defined(__UINTPTR_TYPE__) + +void __ubsan_message(const char *msg); +void __ubsan_message(const char *kind, uintptr_t caller); + +[[noreturn]] void __ubsan_abort(); +[[noreturn]] void __ubsan_abort_with_message(const char *kind, + uintptr_t caller); + +#endif // UBSAN_MINIMAL_COMMON_H diff --git a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp index 6d6638c26ce19..50a3d779ce7e0 100644 --- a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp +++ b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp @@ -1,23 +1,6 @@ +#include "ubsan_minimal_common.h" #include "sanitizer_common/sanitizer_atomic.h" -#include -#include -#include - -#if defined(KERNEL_USE) -extern "C" void ubsan_message(const char *msg); -static void message(const char *msg) { ubsan_message(msg); } -#elif SANITIZER_AMDGPU || SANITIZER_NVPTX -#include -template -static void message(const char *msg, Args &&...args) { - fprintf(stderr, msg, args...); -} -#else -#include -static void message(const char *msg) { (void)write(2, msg, strlen(msg)); } -#endif - // If for some reason we cannot build the runtime with preserve_all, don't // emit any symbol. Programs that need them will fail to link, but that is // better than randomly corrupted registers. @@ -40,45 +23,6 @@ static __sanitizer::atomic_uintptr_t caller_pcs[kMaxCallerPcs]; // that "too many errors" has already been reported. static __sanitizer::atomic_uint32_t caller_pcs_sz; -static char *append_str(const char *s, char *buf, const char *end) { - for (const char *p = s; (buf < end) && (*p != '\0'); ++p, ++buf) - *buf = *p; - return buf; -} - -static char *append_hex(uintptr_t d, char *buf, const char *end) { - // Print the address by nibbles. - for (unsigned shift = sizeof(uintptr_t) * 8; shift && buf < end;) { - shift -= 4; - unsigned nibble = (d >> shift) & 0xf; - *(buf++) = nibble < 10 ? nibble + '0' : nibble - 10 + 'a'; - } - return buf; -} - -static void format_msg(const char *kind, uintptr_t caller, char *buf, - const char *end) { - buf = append_str("ubsan: ", buf, end); - buf = append_str(kind, buf, end); - buf = append_str(" by 0x", buf, end); - buf = append_hex(caller, buf, end); - buf = append_str("\n", buf, end); - if (buf == end) - --buf; // Make sure we don't cause a buffer overflow. - *buf = '\0'; -} - -static void format(const char *kind, uintptr_t caller) { -#if SANITIZER_AMDGPU || SANITIZER_NVPTX - (void)format_msg; - message("ubsan: %s by %p\n", kind, reinterpret_cast(caller)); -#else - char msg_buf[128]; - format_msg(kind, caller, msg_buf, msg_buf + sizeof(msg_buf)); - message(msg_buf); -#endif -} - [[gnu::cold]] static void report_error(const char *kind, uintptr_t caller) { if (caller == 0) return; @@ -106,12 +50,12 @@ static void format(const char *kind, uintptr_t caller) { continue; // Concurrent update! Try again from the start. if (sz == kMaxCallerPcs) { - message("ubsan: too many errors\n"); + __ubsan_message("ubsan: too many errors\n"); return; } __sanitizer::atomic_store_relaxed(&caller_pcs[sz], caller); - format(kind, caller); + __ubsan_message(kind, caller); } } @@ -138,32 +82,19 @@ SANITIZER_INTERFACE_WEAK_DEF(void, __ubsan_report_error_fatal, const char *kind, __ubsan_report_error(kind, caller); } -#if defined(__ANDROID__) -extern "C" __attribute__((weak)) void android_set_abort_message(const char *); static void abort_with_message(const char *kind, uintptr_t caller) { - char msg_buf[128]; - format_msg(kind, caller, msg_buf, msg_buf + sizeof(msg_buf)); - if (&android_set_abort_message) - android_set_abort_message(msg_buf); - abort(); + __ubsan_abort_with_message(kind, caller); } -#elif SANITIZER_AMDGPU || SANITIZER_NVPTX -static void abort_with_message(const char *kind, uintptr_t caller) { - __builtin_verbose_trap("ubsan", "unrecoverable error"); -} -#else -static void abort_with_message(const char *kind, uintptr_t caller) { abort(); } -#endif #if SANITIZER_DEBUG namespace __sanitizer { // The DCHECK macro needs this symbol to be defined. void NORETURN CheckFailed(const char *file, int, const char *cond, u64, u64) { - message("Sanitizer CHECK failed: "); - message(file); - message(":?? : "); // FIXME: Show line number. - message(cond); - abort(); + __ubsan_message("Sanitizer CHECK failed: "); + __ubsan_message(file); + __ubsan_message(":?? : "); // FIXME: Show line number. + __ubsan_message(cond); + __ubsan_abort(); } } // namespace __sanitizer #endif From c1a22693c1f158b8e418a764054e90bf0f3ab8c4 Mon Sep 17 00:00:00 2001 From: Victor Mustya Date: Tue, 5 May 2026 17:52:44 -0700 Subject: [PATCH 2/4] [compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime The SPIR-V target support for UBSan Minimal Runtime is added to provide basic undefined behavior detection capabilities for SPIR-V based devices Assisted-by: Claude Opus 4.6 --- .../cmake/Modules/AllSupportedArchDefs.cmake | 3 +- .../cmake/Modules/CompilerRTUtils.cmake | 3 ++ compiler-rt/cmake/base-config-ix.cmake | 2 + .../lib/sanitizer_common/sanitizer_platform.h | 6 +++ .../lib/ubsan_minimal/targets/spirv64.cpp | 45 +++++++++++++++++++ 5 files changed, 58 insertions(+), 1 deletion(-) create mode 100644 compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp diff --git a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake index fa8b8d6679582..4c9e5c4a81197 100644 --- a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake +++ b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake @@ -15,6 +15,7 @@ set(RISCV64 riscv64) set(S390X s390x) set(SPARC sparc) set(SPARCV9 sparcv9) +set(SPIRV64 spirv64) set(WASM32 wasm32) set(WASM64 wasm64) set(VE ve) @@ -107,7 +108,7 @@ set(ALL_UBSAN_SUPPORTED_ARCH ${X86} ${X86_64} ${ARM32} ${ARM64} ${RISCV64} ${MIPS32} ${MIPS64} ${PPC64} ${S390X} ${SPARC} ${SPARCV9} ${HEXAGON} ${LOONGARCH64}) set(ALL_UBSAN_MINIMAL_SUPPORTED_ARCH - ${ALL_UBSAN_SUPPORTED_ARCH} ${AMDGPU} ${NVPTX}) + ${ALL_UBSAN_SUPPORTED_ARCH} ${AMDGPU} ${NVPTX} ${SPIRV64}) if (OS_NAME MATCHES "FreeBSD") set(ALL_SAFESTACK_SUPPORTED_ARCH ${X86} ${X86_64} ${ARM64}) else() diff --git a/compiler-rt/cmake/Modules/CompilerRTUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTUtils.cmake index dc9d2f4c473f8..2834064b1bede 100644 --- a/compiler-rt/cmake/Modules/CompilerRTUtils.cmake +++ b/compiler-rt/cmake/Modules/CompilerRTUtils.cmake @@ -429,6 +429,9 @@ macro(construct_compiler_rt_default_triple) elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "nvptx|spirv64") set(COMPILER_RT_GPU_BUILD ON) set(CMAKE_REQUIRED_FLAGS "${CMAKE_REQUIRED_FLAGS} -flto -c") + elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "spirv64") + set(COMPILER_RT_GPU_BUILD ON) + set(CMAKE_REQUIRED_FLAGS "${CMAKE_REQUIRED_FLAGS} -flto -c") endif() endif() diff --git a/compiler-rt/cmake/base-config-ix.cmake b/compiler-rt/cmake/base-config-ix.cmake index 4898e3428ebb0..90e6d04daef12 100644 --- a/compiler-rt/cmake/base-config-ix.cmake +++ b/compiler-rt/cmake/base-config-ix.cmake @@ -290,6 +290,8 @@ macro(test_targets) endif() elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "nvptx") test_target_arch(nvptx64 "" "--nvptx64-nvidia-cuda" "-nogpulib" "-flto" "-c") + elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "spirv64") + test_target_arch(spirv64 "" "--target=spirv64-unknown-unknown" "-nogpulib" "-flto" "-c") elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "arm") if(WIN32) test_target_arch(arm "" "" "") diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h index daa62de131fe4..32ffaa4e57c13 100644 --- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h +++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h @@ -318,6 +318,12 @@ # define SANITIZER_NVPTX 0 #endif +#if defined(__SPIRV__) +# define SANITIZER_SPIRV 1 +#else // defined(__SPIRV__) +# define SANITIZER_SPIRV 0 +#endif // defined(__SPIRV__) + // By default we allow to use SizeClassAllocator64 on 64-bit platform. // But in some cases SizeClassAllocator64 does not work well and we need to // fallback to SizeClassAllocator32. diff --git a/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp b/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp new file mode 100644 index 0000000000000..a900d3dfc50ab --- /dev/null +++ b/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp @@ -0,0 +1,45 @@ +#include "ubsan_minimal_common.h" + +using uint3 = unsigned __attribute__((ext_vector_type(3))); + +struct __ubsan_abort_info_t { + __attribute__((opencl_constant)) const char *fmt; + const char *kind; + uintptr_t caller; + uint3 gid; + uint3 lid; +}; + +// OpenCL printf maps to OpExtInst printf (OpenCL extended instruction set). +extern "C" int printf(__attribute__((opencl_constant)) const char *fmt, ...); + +// OpenCL work-item builtins map to SPIR-V BuiltIn variables. +extern "C" unsigned get_group_id(unsigned dim); +extern "C" unsigned get_local_id(unsigned dim); + +static __attribute__((opencl_constant)) const char ubsan_msg_simple[] = "%s"; +static __attribute__((opencl_constant)) const char ubsan_msg_fmt[] = + "ubsan: %s by 0x%lx at gid=[%v3u] lid=[%v3u]\n"; + +void __ubsan_message(const char *msg) { printf(ubsan_msg_simple, msg); } + +void __ubsan_message(const char *kind, uintptr_t caller) { + uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; + uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; + + printf(ubsan_msg_fmt, kind, caller, gid, lid); +} + +// SPV_KHR_abort: OpAbortKHR terminates the invocation and passes a message +// to the client API. The message is passed as a typed value. +[[noreturn]] void __spirv_AbortKHR(__ubsan_abort_info_t info); + +[[noreturn]] void __ubsan_abort() { __ubsan_abort_with_message("abort", 0); } + +[[noreturn]] void __ubsan_abort_with_message(const char *kind, + uintptr_t caller) { + uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; + uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; + __ubsan_abort_info_t info = {ubsan_msg_fmt, kind, caller, gid, lid}; + __spirv_AbortKHR(info); +} From 06ffdd645b691ea427baa13198d61f2b98243561 Mon Sep 17 00:00:00 2001 From: Victor Mustya Date: Wed, 6 May 2026 08:45:23 -0700 Subject: [PATCH 3/4] Fix inconsistencies in CMake and the header --- compiler-rt/cmake/Modules/CompilerRTUtils.cmake | 3 --- compiler-rt/cmake/config-ix.cmake | 6 ++---- compiler-rt/lib/sanitizer_common/sanitizer_platform.h | 4 ++-- 3 files changed, 4 insertions(+), 9 deletions(-) diff --git a/compiler-rt/cmake/Modules/CompilerRTUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTUtils.cmake index 2834064b1bede..dc9d2f4c473f8 100644 --- a/compiler-rt/cmake/Modules/CompilerRTUtils.cmake +++ b/compiler-rt/cmake/Modules/CompilerRTUtils.cmake @@ -429,9 +429,6 @@ macro(construct_compiler_rt_default_triple) elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "nvptx|spirv64") set(COMPILER_RT_GPU_BUILD ON) set(CMAKE_REQUIRED_FLAGS "${CMAKE_REQUIRED_FLAGS} -flto -c") - elseif("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "spirv64") - set(COMPILER_RT_GPU_BUILD ON) - set(CMAKE_REQUIRED_FLAGS "${CMAKE_REQUIRED_FLAGS} -flto -c") endif() endif() diff --git a/compiler-rt/cmake/config-ix.cmake b/compiler-rt/cmake/config-ix.cmake index 8129fb662c72d..008086f598acc 100644 --- a/compiler-rt/cmake/config-ix.cmake +++ b/compiler-rt/cmake/config-ix.cmake @@ -893,11 +893,9 @@ else() set(COMPILER_RT_HAS_UBSAN FALSE) endif() -if (UBSAN_SUPPORTED_ARCH AND +if (UBSAN_MINIMAL_SUPPORTED_ARCH AND (OS_NAME MATCHES "Linux|FreeBSD|NetBSD|Android|Darwin|SunOS" OR - "${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn|nvptx")) - set(COMPILER_RT_HAS_UBSAN_MINIMAL TRUE) -elseif(COMPILER_RT_GPU_BUILD AND UBSAN_MINIMAL_SUPPORTED_ARCH) + COMPILER_RT_GPU_BUILD)) set(COMPILER_RT_HAS_UBSAN_MINIMAL TRUE) else() set(COMPILER_RT_HAS_UBSAN_MINIMAL FALSE) diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h index 32ffaa4e57c13..5bcd3772486bc 100644 --- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h +++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h @@ -320,9 +320,9 @@ #if defined(__SPIRV__) # define SANITIZER_SPIRV 1 -#else // defined(__SPIRV__) +#else # define SANITIZER_SPIRV 0 -#endif // defined(__SPIRV__) +#endif // By default we allow to use SizeClassAllocator64 on 64-bit platform. // But in some cases SizeClassAllocator64 does not work well and we need to From 85544438d337bff1a50026b36db51f60fa0ba081 Mon Sep 17 00:00:00 2001 From: Victor Mustya Date: Wed, 6 May 2026 09:43:35 -0700 Subject: [PATCH 4/4] Use Clang SPIR-V built-in functions, when it's possible --- .../lib/ubsan_minimal/targets/spirv64.cpp | 44 +++++++++++-------- .../lib/ubsan_minimal/ubsan_minimal_common.h | 7 +-- 2 files changed, 29 insertions(+), 22 deletions(-) diff --git a/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp b/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp index a900d3dfc50ab..52cbcb32658a1 100644 --- a/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp +++ b/compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp @@ -1,45 +1,51 @@ #include "ubsan_minimal_common.h" -using uint3 = unsigned __attribute__((ext_vector_type(3))); +#define __gpu_constant __attribute__((opencl_constant)) +using size_t3 = size_t [[clang::ext_vector_type(3)]]; struct __ubsan_abort_info_t { - __attribute__((opencl_constant)) const char *fmt; + __gpu_constant const char *fmt; const char *kind; uintptr_t caller; - uint3 gid; - uint3 lid; + size_t3 gid; + size_t3 lid; }; // OpenCL printf maps to OpExtInst printf (OpenCL extended instruction set). -extern "C" int printf(__attribute__((opencl_constant)) const char *fmt, ...); +extern "C" int printf(__gpu_constant const char *fmt, ...); -// OpenCL work-item builtins map to SPIR-V BuiltIn variables. -extern "C" unsigned get_group_id(unsigned dim); -extern "C" unsigned get_local_id(unsigned dim); +// SPV_KHR_abort: OpAbortKHR terminates the invocation and passes a message +// to the client API. The message is passed as a typed value. +[[noreturn]] void __spirv_AbortKHR(__ubsan_abort_info_t info); -static __attribute__((opencl_constant)) const char ubsan_msg_simple[] = "%s"; -static __attribute__((opencl_constant)) const char ubsan_msg_fmt[] = - "ubsan: %s by 0x%lx at gid=[%v3u] lid=[%v3u]\n"; +static __gpu_constant const char ubsan_msg_simple[] = "%s"; +static __gpu_constant const char ubsan_msg_fmt[] = + "ubsan: %s by 0x%lx at gid=[%v3lu] lid=[%v3lu]\n"; void __ubsan_message(const char *msg) { printf(ubsan_msg_simple, msg); } void __ubsan_message(const char *kind, uintptr_t caller) { - uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; - uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; + size_t3 gid = {__builtin_spirv_workgroup_id(0), + __builtin_spirv_workgroup_id(1), + __builtin_spirv_workgroup_id(2)}; + size_t3 lid = {__builtin_spirv_local_invocation_id(0), + __builtin_spirv_local_invocation_id(1), + __builtin_spirv_local_invocation_id(2)}; printf(ubsan_msg_fmt, kind, caller, gid, lid); } -// SPV_KHR_abort: OpAbortKHR terminates the invocation and passes a message -// to the client API. The message is passed as a typed value. -[[noreturn]] void __spirv_AbortKHR(__ubsan_abort_info_t info); - [[noreturn]] void __ubsan_abort() { __ubsan_abort_with_message("abort", 0); } [[noreturn]] void __ubsan_abort_with_message(const char *kind, uintptr_t caller) { - uint3 gid = {get_group_id(0), get_group_id(1), get_group_id(2)}; - uint3 lid = {get_local_id(0), get_local_id(1), get_local_id(2)}; + size_t3 gid = {__builtin_spirv_workgroup_id(0), + __builtin_spirv_workgroup_id(1), + __builtin_spirv_workgroup_id(2)}; + size_t3 lid = {__builtin_spirv_local_invocation_id(0), + __builtin_spirv_local_invocation_id(1), + __builtin_spirv_local_invocation_id(2)}; + __ubsan_abort_info_t info = {ubsan_msg_fmt, kind, caller, gid, lid}; __spirv_AbortKHR(info); } diff --git a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h index 0a665df7e50c3..c17bf1049eb99 100644 --- a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h +++ b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h @@ -13,14 +13,15 @@ #ifndef UBSAN_MINIMAL_COMMON_H #define UBSAN_MINIMAL_COMMON_H -#if defined(__UINTPTR_TYPE__) +#if defined(__UINTPTR_TYPE__) && defined(__SIZE_TYPE__) typedef __UINTPTR_TYPE__ uintptr_t; +typedef __SIZE_TYPE__ size_t; static_assert(sizeof(uintptr_t) == sizeof(void *), "uintptr_t must be the same size as void*"); -#else // defined(__UINTPTR_TYPE__) +#else // defined(__UINTPTR_TYPE__) && defined(__SIZE_TYPE__) #include -#endif // defined(__UINTPTR_TYPE__) +#endif // defined(__UINTPTR_TYPE__) && defined(__SIZE_TYPE__) void __ubsan_message(const char *msg); void __ubsan_message(const char *kind, uintptr_t caller);