[compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime#195979
Closed
vmustya wants to merge 4 commits into
Closed
[compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime#195979vmustya wants to merge 4 commits into
vmustya wants to merge 4 commits into
Conversation
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.
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 <noreply@anthropic.com>
llvmorg-github-actions
Bot
added
compiler-rt
compiler-rt:ubsan
|
@llvm/pr-subscribers-compiler-rt-sanitizer Author: Victor Mustya (vmustya) ChangesThe SPIR-V target support for UBSan Minimal Runtime is added to provide Assisted-by: Claude Opus 4.6 <noreply@anthropic.com> Full diff: https://github.com/llvm/llvm-project/pull/195979.diff 11 Files Affected:
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..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)
@@ -105,7 +106,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} ${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/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..32ffaa4e57c13 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
@@ -314,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/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 $<IF:$<BOOL:${COMPILER_RT_GPU_BUILD}>,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 <stdlib.h>
+#include <string.h>
+
+#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 <stdio.h>
+static void message(const char *msg) { fprintf(stderr, "%s", msg); }
+#else
+#include <unistd.h>
+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/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);
+}
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 <stdint.h>
+#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 <stdint.h>
-#include <stdlib.h>
-#include <string.h>
-
-#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 <stdio.h>
-template <typename... Args>
-static void message(const char *msg, Args &&...args) {
- fprintf(stderr, msg, args...);
-}
-#else
-#include <unistd.h>
-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<void *>(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
|
You can test this locally with the following command:git-clang-format --diff origin/main HEAD --extensions h,cpp -- compiler-rt/lib/ubsan_minimal/targets/generic.cpp compiler-rt/lib/ubsan_minimal/targets/spirv64.cpp compiler-rt/lib/ubsan_minimal/ubsan_minimal_common.h compiler-rt/lib/sanitizer_common/sanitizer_platform.h compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp --diff_from_common_commit
View the diff from clang-format here.diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
index 5bcd37724..63b8a772f 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
@@ -14,7 +14,7 @@
#if !defined(SANITIZER_GPU)
# define SANITIZER_GPU 0
-#endif // !defined(SANITIZER_GPU)
+#endif // !defined(SANITIZER_GPU)
#if !defined(__linux__) && !defined(__FreeBSD__) && !defined(__NetBSD__) && \
!defined(__APPLE__) && !defined(_WIN32) && !defined(__Fuchsia__) && \
diff --git a/compiler-rt/lib/ubsan_minimal/targets/generic.cpp b/compiler-rt/lib/ubsan_minimal/targets/generic.cpp
index 41c340a35..17911d7a9 100644
--- a/compiler-rt/lib/ubsan_minimal/targets/generic.cpp
+++ b/compiler-rt/lib/ubsan_minimal/targets/generic.cpp
@@ -47,9 +47,7 @@ static void format_msg(const char *kind, uintptr_t caller, char *buf,
*buf = '\0';
}
-void __ubsan_message(const char *msg) {
- message(msg);
-}
+void __ubsan_message(const char *msg) { message(msg); }
void __ubsan_message(const char *kind, uintptr_t caller) {
char buf[128];
@@ -57,9 +55,7 @@ void __ubsan_message(const char *kind, uintptr_t caller) {
message(buf);
}
-void __ubsan_abort() {
- abort();
-}
+void __ubsan_abort() { abort(); }
void __ubsan_abort_with_message(const char *kind, uintptr_t caller) {
char buf[128];
@@ -68,7 +64,7 @@ void __ubsan_abort_with_message(const char *kind, uintptr_t caller) {
#if defined(__ANDROID__)
if (&android_set_abort_message)
android_set_abort_message(buf);
-#else // defined(__ANDROID__)
+#else // defined(__ANDROID__)
message(buf);
#endif // defined(__ANDROID__)
diff --git a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp
index 50a3d779c..4c9a154bd 100644
--- a/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp
+++ b/compiler-rt/lib/ubsan_minimal/ubsan_minimal_handlers.cpp
@@ -1,5 +1,5 @@
-#include "ubsan_minimal_common.h"
#include "sanitizer_common/sanitizer_atomic.h"
+#include "ubsan_minimal_common.h"
// 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
|
Contributor
Author
|
This PR depends on the #193037, #195976 and #195978. @jhuber6, @vitalybuka, could you please review? Is it okay to split the minimal runtime to multiple source files instead of having everything in a single one under ifdef? Thanks in advance! |
jhuber6
requested changes
May 6, 2026
jhuber6
left a comment
jhuber6
left a comment
Contributor
There was a problem hiding this comment.
Any reason we need to explicitly match on spirv64 and not spirv?
| 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") |
Contributor
There was a problem hiding this comment.
Isn't this case handled above?
Comment on lines
+900
to
+901
| elseif(COMPILER_RT_GPU_BUILD AND UBSAN_MINIMAL_SUPPORTED_ARCH) | ||
| set(COMPILER_RT_HAS_UBSAN_MINIMAL TRUE) |
| 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}) |
|
|
||
| #if defined(__SPIRV__) | ||
| # define SANITIZER_SPIRV 1 | ||
| #else // defined(__SPIRV__) |
|
|
||
| // 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); |
Contributor
There was a problem hiding this comment.
This should not be defined here, it should be what __builtin_verbose_trap or __builtin_trap lowers to.
Contributor
Author
There was a problem hiding this comment.
The __builtin_verbose_trap doesn't produce any semantic info in the SPIR-V (or any other generated code). It generates the llvm.trap intrinsic and some extra debug info. However, the __spirv_AbortKHR is lowered into the OpAbortKHR SPIR-V instruction with an arbitrary error information. Here the sanitizer message is passed as the abort operand, so the offload device compiler would be able to print it before the kernel termination.
See the SPIR-V extension spec for more info: https://github.khronos.org/SPIRV-Registry/extensions/KHR/SPV_KHR_abort.html
Contributor
There was a problem hiding this comment.
Trapping isn't the intended way for the minimal runtime to work. If you just want to abort on first error you'd use -fsanitize-trap=all. There's not much value in having a minimal runtime that functions the same as the trapping version.
Contributor
Author
There was a problem hiding this comment.
The idea is to report an error and continue, when the error is recoverable. There are two kinds of the error handlers: non-aborting and aborting. For the non-aborting handlers we just should the printf and continue.
Also, the -fsanitize-trap=all only emits the trap instruction, without any error message. It's hard to work with such traps, when the debugger support is limited.
| @@ -0,0 +1,45 @@ | |||
| #include "ubsan_minimal_common.h" | |||
|
|
|||
| using uint3 = unsigned __attribute__((ext_vector_type(3))); | |||
Contributor
There was a problem hiding this comment.
Why is all of this necessary?
Contributor
Author
There was a problem hiding this comment.
I believe, that the group ID and local ID info of the error should be useful for the GPU targets. Here is the OpenCL printf used. It supports vector printing -- the %v3u in the format string. Such parameters should be passed as vectors.
Reference: https://registry.khronos.org/OpenCL/specs/unified/html/OpenCL_C.html#printf-format-string
| @@ -0,0 +1,76 @@ | |||
|
|
|||
| #include "ubsan_minimal_common.h" | |||
Contributor
There was a problem hiding this comment.
Not a fan of splitting this up so arbitrarily, I don't know what SPIR-V does that can't be done by improving the set of compiler builtins
Contributor
Author
There was a problem hiding this comment.
The __spirv_* functions are actually the compiler built-ins for the SPIRV target. So, I'm not sure I understand how we could improve the built-ins.
Contributor
There was a problem hiding this comment.
I guess I don't really understand how the SPIR-V build is supposed to work, but why do we have both C libcalls and stuff like __builtin_spirv_num_workgroups(0)? I generally prefer the latter.
Contributor
Author
There was a problem hiding this comment.
I've updated the file to use the __builtin_spirv_* functions when it's possible. Still, the __spirv_AbortKHR(...) is used, because it's handled by the backend. 8554443
jhuber6
reviewed
May 6, 2026
jhuber6
left a comment
jhuber6
left a comment
Contributor
There was a problem hiding this comment.
What is the main delta between the old code? Like like to avoid this new indirection if possible. Both should just have access to printf, right? If you want to print the grid dimensions that's just some extra arguments to the existing case. I'm personally not sure we gain too much from knowing which thread hits the UB
Contributor
Author
The existing code tries to format the error message on a stack buffer and then print it. That's what we can't support for GPU in common case. The SPIR-V target, the Also, the old code uses the |
vitalybuka
requested changes
May 6, 2026
Contributor
There was a problem hiding this comment.
I suggest keep this PR open and extract small refactorings into own PRs
https://releases.llvm.org/14.0.0/docs/DeveloperPolicy.html#id19
if they are just opportunistic unrelated improvements, state it there (or consider abandoning).
if they are realy needed for SPIR-V PR, please still consider extracting into own PRs, just refer this one in description and why it's needed
Contributor
That's news to me, it's supposed to just call printf. Does this not work in SPIR-V? Their printf support isn't very useful if not. 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<void *>(caller));
#else
char msg_buf[128];
format_msg(kind, caller, msg_buf, msg_buf + sizeof(msg_buf));
message(msg_buf);
#endif
}
Are you reading a really old version of this? That calls |
Contributor
Author
|
@jhuber6, @vitalybuka, thank you for your feedback. I really appreciate it. I've looked at the latest compiler-rt state more carefully, and I've decided to implement the SPIR-V support without any file splitting so far. I've created a new PR #196149. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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 noreply@anthropic.com