[compiler-rt] [compiler-rt] Add SPIR-V target support for UBSan Minimal Runtime (PR #195979)
Victor Mustya via llvm-commits
llvm-commits at lists.llvm.org
Wed May 6 09:44:23 PDT 2026
https://github.com/vmustya updated https://github.com/llvm/llvm-project/pull/195979
>From eab515b6d9519dfab6866aab0185b51c183a0390 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
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 $<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/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
>From c1a22693c1f158b8e418a764054e90bf0f3ab8c4 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
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 <noreply at anthropic.com>
---
.../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 <victor.mustya at intel.com>
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 <victor.mustya at intel.com>
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 <stdint.h>
-#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);
More information about the llvm-commits
mailing list