[llvm-branch-commits] [compiler-rt] [compiler-rt] Add 'csan' library for the concurrency sanitizer (PR #225782)
Joseph Huber via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Oct 8 13:52:35 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/225782
>From 79e313e6dbde8cba70a638f8723ed3f0bd440ad9 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Mon, 21 Sep 2026 22:14:12 -0500
Subject: [PATCH 1/5] [compiler-rt] Add 'csan' library for the concurrency
sanitizer
Summary:
Adds the runtime for the concurrency sanitizer, both CPU and GPU.
Fundamentally, this works using the following pseudocode:
```c
static u64 watchpoints[N]; // Hash-indexed, zero is empty.
// Emitted before the access, so we never trip on our own write.
void check_access(volatile void *addr, u32 size, u32 type) {
// Every access probes. A read conflicts only with a watched write, a
// write conflicts with either.
if (u64 *wp = find_watchpoint(addr, size, type))
consume(wp, this_pc()); // Hand our location to the owner.
if (!should_sample()) // Wave-uniform, 1-in-N chance.
return;
u64 *wp = arm_watchpoint(addr, size, type);
if (!wp) // Slot is already taken.
return;
auto old = read(addr, size);
delay(ctx.rand()); // Randomized, bounded.
if (void *peer = disarm(wp))
report_race(addr, this_pc(), peer);
else if (read(addr, size) != old)
report_race(addr, this_pc(), UNKNOWN);
}
```
This implementation is kept intentionally minimal to simplify the review
process. Many options are planned for later. The intended use is for
users to pass `-fsanitize=concurrency` for supported compilations.
---
compiler-rt/cmake/caches/AMDGPU.cmake | 2 +-
compiler-rt/cmake/config-ix.cmake | 11 +-
compiler-rt/lib/csan/CMakeLists.txt | 60 +++
compiler-rt/lib/csan/csan.cpp | 334 +++++++++++++
compiler-rt/lib/csan/csan.h | 64 +++
compiler-rt/lib/csan/csan_defs.h | 35 ++
compiler-rt/lib/csan/csan_flags.inc | 23 +
compiler-rt/lib/csan/csan_gpu.cpp | 456 ++++++++++++++++++
compiler-rt/lib/csan/csan_offload_packet.h | 39 ++
compiler-rt/lib/csan/csan_report.cpp | 279 +++++++++++
compiler-rt/lib/csan/csan_watch.h | 163 +++++++
compiler-rt/lib/csan/offload/CMakeLists.txt | 25 +
compiler-rt/lib/csan/offload/csan_offload.h | 31 ++
.../offload/csan_offload_hsa_interceptors.cpp | 370 ++++++++++++++
.../lib/csan/offload/csan_offload_report.cpp | 184 +++++++
.../sanitizer_internal_defs.h | 3 +
.../sanitizer_offload_opcodes.h | 1 +
compiler-rt/test/csan/AMDGPU/aba-race.hip | 33 ++
compiler-rt/test/csan/AMDGPU/array-race.hip | 22 +
.../test/csan/AMDGPU/atomic-nonatomic.hip | 26 +
compiler-rt/test/csan/AMDGPU/atomic.hip | 17 +
compiler-rt/test/csan/AMDGPU/disjoint.hip | 20 +
compiler-rt/test/csan/AMDGPU/global-race.hip | 23 +
compiler-rt/test/csan/AMDGPU/helper-race.hip | 26 +
compiler-rt/test/csan/AMDGPU/lds-disjoint.hip | 21 +
compiler-rt/test/csan/AMDGPU/lds-race.hip | 21 +
compiler-rt/test/csan/AMDGPU/lit.local.cfg.py | 11 +
.../test/csan/AMDGPU/memcpy-large-race.hip | 28 ++
compiler-rt/test/csan/AMDGPU/memcpy-race.hip | 22 +
compiler-rt/test/csan/AMDGPU/memmove-race.hip | 28 ++
compiler-rt/test/csan/AMDGPU/openmp-race.cpp | 19 +
compiler-rt/test/csan/AMDGPU/race.h | 32 ++
.../test/csan/AMDGPU/shared-watchpoints.hip | 59 +++
.../test/csan/AMDGPU/single-thread.hip | 17 +
.../test/csan/AMDGPU/write-read-race.hip | 26 +
compiler-rt/test/csan/CMakeLists.txt | 25 +
compiler-rt/test/csan/access-sizes.cpp | 62 +++
compiler-rt/test/csan/atomic-nonatomic.cpp | 27 ++
compiler-rt/test/csan/atomic.cpp | 22 +
compiler-rt/test/csan/halt-on-error.cpp | 24 +
compiler-rt/test/csan/ignore-thread.cpp | 45 ++
compiler-rt/test/csan/large-access.cpp | 19 +
compiler-rt/test/csan/large-range-race.cpp | 25 +
compiler-rt/test/csan/lit.common.cfg.py | 62 +++
compiler-rt/test/csan/lit.site.cfg.py.in | 8 +
compiler-rt/test/csan/race.cpp | 26 +
compiler-rt/test/csan/read-read.cpp | 23 +
compiler-rt/test/csan/single-thread.cpp | 11 +
compiler-rt/test/csan/unknown-origin.cpp | 31 ++
compiler-rt/test/lit.common.cfg.py | 1 +
50 files changed, 2940 insertions(+), 2 deletions(-)
create mode 100644 compiler-rt/lib/csan/CMakeLists.txt
create mode 100644 compiler-rt/lib/csan/csan.cpp
create mode 100644 compiler-rt/lib/csan/csan.h
create mode 100644 compiler-rt/lib/csan/csan_defs.h
create mode 100644 compiler-rt/lib/csan/csan_flags.inc
create mode 100644 compiler-rt/lib/csan/csan_gpu.cpp
create mode 100644 compiler-rt/lib/csan/csan_offload_packet.h
create mode 100644 compiler-rt/lib/csan/csan_report.cpp
create mode 100644 compiler-rt/lib/csan/csan_watch.h
create mode 100644 compiler-rt/lib/csan/offload/CMakeLists.txt
create mode 100644 compiler-rt/lib/csan/offload/csan_offload.h
create mode 100644 compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
create mode 100644 compiler-rt/lib/csan/offload/csan_offload_report.cpp
create mode 100644 compiler-rt/test/csan/AMDGPU/aba-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/array-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/atomic-nonatomic.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/atomic.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/disjoint.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/global-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/helper-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/lds-disjoint.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/lds-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/lit.local.cfg.py
create mode 100644 compiler-rt/test/csan/AMDGPU/memcpy-large-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/memcpy-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/memmove-race.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/openmp-race.cpp
create mode 100644 compiler-rt/test/csan/AMDGPU/race.h
create mode 100644 compiler-rt/test/csan/AMDGPU/shared-watchpoints.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/single-thread.hip
create mode 100644 compiler-rt/test/csan/AMDGPU/write-read-race.hip
create mode 100644 compiler-rt/test/csan/CMakeLists.txt
create mode 100644 compiler-rt/test/csan/access-sizes.cpp
create mode 100644 compiler-rt/test/csan/atomic-nonatomic.cpp
create mode 100644 compiler-rt/test/csan/atomic.cpp
create mode 100644 compiler-rt/test/csan/halt-on-error.cpp
create mode 100644 compiler-rt/test/csan/ignore-thread.cpp
create mode 100644 compiler-rt/test/csan/large-access.cpp
create mode 100644 compiler-rt/test/csan/large-range-race.cpp
create mode 100644 compiler-rt/test/csan/lit.common.cfg.py
create mode 100644 compiler-rt/test/csan/lit.site.cfg.py.in
create mode 100644 compiler-rt/test/csan/race.cpp
create mode 100644 compiler-rt/test/csan/read-read.cpp
create mode 100644 compiler-rt/test/csan/single-thread.cpp
create mode 100644 compiler-rt/test/csan/unknown-origin.cpp
diff --git a/compiler-rt/cmake/caches/AMDGPU.cmake b/compiler-rt/cmake/caches/AMDGPU.cmake
index bb5ab22edaf69e..34a5559fe52990 100644
--- a/compiler-rt/cmake/caches/AMDGPU.cmake
+++ b/compiler-rt/cmake/caches/AMDGPU.cmake
@@ -7,7 +7,7 @@ set(COMPILER_RT_BUILD_BUILTINS ON CACHE BOOL "")
set(COMPILER_RT_BAREMETAL_BUILD ON CACHE BOOL "")
set(COMPILER_RT_BUILD_CRT OFF CACHE BOOL "")
set(COMPILER_RT_BUILD_SANITIZERS ON CACHE BOOL "")
-set(COMPILER_RT_SANITIZERS_TO_BUILD "ubsan;ubsan_minimal" CACHE STRING "")
+set(COMPILER_RT_SANITIZERS_TO_BUILD "ubsan;ubsan_minimal;csan" CACHE STRING "")
set(COMPILER_RT_BUILD_XRAY OFF CACHE BOOL "")
set(COMPILER_RT_BUILD_LIBFUZZER OFF CACHE BOOL "")
set(COMPILER_RT_BUILD_PROFILE ON CACHE BOOL "")
diff --git a/compiler-rt/cmake/config-ix.cmake b/compiler-rt/cmake/config-ix.cmake
index d411cf014dd34b..d1732e761e119b 100644
--- a/compiler-rt/cmake/config-ix.cmake
+++ b/compiler-rt/cmake/config-ix.cmake
@@ -768,7 +768,7 @@ if(COMPILER_RT_SUPPORTED_ARCH)
endif()
message(STATUS "Compiler-RT supported architectures: ${COMPILER_RT_SUPPORTED_ARCH}")
-set(ALL_SANITIZERS asan;rtsan;dfsan;msan;hwasan;tsan;tysan;safestack;cfi;scudo_standalone;ubsan_minimal;gwp_asan;nsan;asan_abi)
+set(ALL_SANITIZERS asan;rtsan;dfsan;msan;hwasan;tsan;csan;tysan;safestack;cfi;scudo_standalone;ubsan_minimal;gwp_asan;nsan;asan_abi)
set(COMPILER_RT_SANITIZERS_TO_BUILD all CACHE STRING
"sanitizers to build if supported on the target (all;${ALL_SANITIZERS})")
list_replace(COMPILER_RT_SANITIZERS_TO_BUILD all "${ALL_SANITIZERS}")
@@ -897,6 +897,15 @@ else()
set(COMPILER_RT_HAS_UBSAN FALSE)
endif()
+if ((COMPILER_RT_HAS_SANITIZER_COMMON AND UBSAN_SUPPORTED_ARCH AND
+ OS_NAME MATCHES "Linux" AND NOT ANDROID)
+ OR (COMPILER_RT_GPU_BUILD AND COMPILER_RT_TARGET_AMDGPU AND
+ UBSAN_SUPPORTED_ARCH))
+ set(COMPILER_RT_HAS_CSAN TRUE)
+else()
+ set(COMPILER_RT_HAS_CSAN FALSE)
+endif()
+
if (UBSAN_SUPPORTED_ARCH AND
(OS_NAME MATCHES "Linux|FreeBSD|NetBSD|Android|Darwin|SunOS" OR
COMPILER_RT_GPU_BUILD))
diff --git a/compiler-rt/lib/csan/CMakeLists.txt b/compiler-rt/lib/csan/CMakeLists.txt
new file mode 100644
index 00000000000000..b0791c39e2a4cf
--- /dev/null
+++ b/compiler-rt/lib/csan/CMakeLists.txt
@@ -0,0 +1,60 @@
+# Build for the ConcurrencySanitizer runtime support library.
+
+include_directories(.)
+include_directories(..)
+include_directories(../../include)
+
+# GPU targets only need the watchpoint runtime. The host interceptor archive is
+# built in the offload/ project.
+if(COMPILER_RT_GPU_BUILD)
+ include(FindLibcCommonUtils)
+ if(NOT TARGET llvm-libc-common-utilities)
+ add_compiler_rt_component(csan)
+ return()
+ endif()
+
+ set(CSAN_SOURCES csan_gpu.cpp)
+ set(CSAN_HEADERS csan_defs.h csan_offload_packet.h csan_watch.h)
+
+ set(CSAN_CFLAGS
+ ${SANITIZER_COMMON_CFLAGS}
+ -DSANITIZER_COMMON_NO_REDEFINE_BUILTINS)
+ append_rtti_flag(OFF CSAN_CFLAGS)
+
+ add_compiler_rt_component(csan)
+
+ add_compiler_rt_runtime(clang_rt.csan
+ STATIC
+ ARCHS ${UBSAN_SUPPORTED_ARCH}
+ SOURCES ${CSAN_SOURCES}
+ ADDITIONAL_HEADERS ${CSAN_HEADERS}
+ CFLAGS ${CSAN_CFLAGS}
+ LINK_LIBS llvm-libc-common-utilities
+ PARENT_TARGET csan)
+
+ return()
+endif()
+
+set(CSAN_CFLAGS ${SANITIZER_COMMON_CFLAGS})
+append_rtti_flag(OFF CSAN_CFLAGS)
+
+add_compiler_rt_component(csan)
+
+add_compiler_rt_runtime(clang_rt.csan
+ STATIC
+ ARCHS ${UBSAN_SUPPORTED_ARCH}
+ SOURCES csan.cpp csan_report.cpp
+ ADDITIONAL_HEADERS csan.h csan_defs.h csan_flags.inc csan_watch.h
+ OBJECT_LIBS RTSanitizerCommon
+ RTSanitizerCommonLibc
+ RTSanitizerCommonCoverage
+ RTSanitizerCommonSymbolizer
+ RTSanitizerCommonSymbolizerInternal
+ RTInterception
+ CFLAGS ${CSAN_CFLAGS}
+ PARENT_TARGET csan)
+
+# Thin host interceptors for offloading. Linked with -u to keep it separate.
+if(OS_NAME MATCHES "Linux" AND NOT ANDROID)
+ add_subdirectory(offload)
+endif()
diff --git a/compiler-rt/lib/csan/csan.cpp b/compiler-rt/lib/csan/csan.cpp
new file mode 100644
index 00000000000000..148e30643333ae
--- /dev/null
+++ b/compiler-rt/lib/csan/csan.cpp
@@ -0,0 +1,334 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Watchpoint-based host data race detector inspired by KCSAN. The host and GPU
+/// runtimes share a configurable packed-watchpoint implementation.
+///
+//===----------------------------------------------------------------------===//
+
+#include "csan.h"
+#include "csan_watch.h"
+
+#include "sanitizer_common/sanitizer_atomic.h"
+#include "sanitizer_common/sanitizer_common.h"
+#include "sanitizer_common/sanitizer_flag_parser.h"
+#include "sanitizer_common/sanitizer_flags.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+#include "sanitizer_common/sanitizer_libc.h"
+#include "sanitizer_common/sanitizer_mutex.h"
+#include "sanitizer_common/sanitizer_stacktrace.h"
+#include "sanitizer_common/sanitizer_symbolizer.h"
+
+using namespace __sanitizer;
+
+#define INTERFACE extern "C" SANITIZER_INTERFACE_ATTRIBUTE
+
+static constexpr u32 kHostWatchpointEntries = 64;
+static constexpr u32 kHostCheckAdjacent = 1;
+static constexpr uptr kHostSlotRange = 4096;
+static constexpr uptr kHostMaxAccessSize =
+ kHostSlotRange * (1 + kHostCheckAdjacent);
+static_assert(__atomic_always_lock_free(sizeof(u64), nullptr),
+ "host watchpoints must be lock-free");
+
+using HostWatchpointTable =
+ __csan::WatchpointTable<kHostMaxAccessSize, kHostCheckAdjacent>;
+static u64 HostWatchpoints[kHostWatchpointEntries +
+ HostWatchpointTable::OverflowEntries];
+static_assert(HostWatchpointTable::AddressBits == 48,
+ "host watchpoints require 48-bit pointers");
+
+static HostWatchpointTable GetHostWatchpoints() {
+ return HostWatchpointTable(HostWatchpoints);
+}
+
+static u32 HostWatchpointSlot(uptr Address) {
+ return Address / kHostSlotRange % kHostWatchpointEntries;
+}
+
+namespace __csan {
+
+Flags flags_data;
+
+static void Initialize();
+
+static atomic_uint64_t NumDataRaces;
+static THREADLOCAL u32 DisableCount;
+static THREADLOCAL s32 Skip;
+static THREADLOCAL u32 RandState;
+
+namespace {
+struct ScopedDisable {
+ ScopedDisable() { ++DisableCount; }
+ ~ScopedDisable() { --DisableCount; }
+};
+} // namespace
+
+void RecordDataRace() {
+ atomic_fetch_add(&NumDataRaces, 1, memory_order_relaxed);
+}
+
+void Flags::SetDefaults() {
+#define CSAN_FLAG(Type, Name, DefaultValue, Description) Name = DefaultValue;
+#include "csan_flags.inc"
+#undef CSAN_FLAG
+}
+
+static void RegisterCsanFlags(FlagParser *Parser, Flags *F) {
+#define CSAN_FLAG(Type, Name, DefaultValue, Description) \
+ RegisterFlag(Parser, #Name, Description, &F->Name);
+#include "csan_flags.inc"
+#undef CSAN_FLAG
+}
+
+void InitializeFlags() {
+ SetCommonFlagsDefaults();
+ {
+ CommonFlags CF;
+ CF.CopyFrom(*common_flags());
+ CF.external_symbolizer_path = GetEnv("CSAN_SYMBOLIZER_PATH");
+ OverrideCommonFlags(CF);
+ }
+
+ flags()->SetDefaults();
+
+ FlagParser Parser;
+ RegisterCommonFlags(&Parser);
+ RegisterCsanFlags(&Parser, flags());
+ Parser.ParseString(__csan_default_options());
+ Parser.ParseStringFromEnv("CSAN_OPTIONS");
+ InitializeCommonFlags();
+ if (Verbosity())
+ ReportUnrecognizedFlags();
+ if (common_flags()->help)
+ Parser.PrintFlagDescriptions();
+}
+
+static u32 Random(u32 EpRo) {
+ if (EpRo <= 1)
+ return 0;
+ u32 State = RandState;
+ if (!State)
+ State = (u32)__builtin_readcyclecounter() | 1u;
+ State = 1664525u * State + 1013904223u;
+ RandState = State;
+ return State % EpRo;
+}
+
+static void ResetSkip() {
+ s32 Count = flags()->skip_watch;
+ if (Count < 0)
+ Count = 0;
+ if (Count)
+ Count -= (s32)Random((u32)Count);
+ Skip = Count;
+}
+
+static bool ShouldWatch(int Type) {
+ if (Type & CSAN_ACCESS_ATOMIC)
+ return false;
+ if (--Skip >= 0)
+ return false;
+ return true;
+}
+
+static void DelayAccess(int Type) {
+ s32 Delay = flags()->udelay;
+ if (Delay < 0)
+ Delay = 0;
+ if (Delay) {
+ u32 Skew = (Type & CSAN_ACCESS_COMPOUND) ? 1u : 0u;
+ u32 Span = (u32)Delay >> Skew;
+ if (!Span)
+ Span = (u32)Delay;
+ Delay -= (s32)Random(Span);
+ }
+ if (Delay)
+ internal_usleep((u64)Delay);
+}
+
+static u64 ReadRange(const volatile u8 *Bytes, uptr Size) {
+ u64 Sum = 0xcbf29ce484222325ull;
+ uptr I = 0;
+ for (; I < Size && ((uptr)(Bytes + I) & 7u); ++I)
+ Sum = (Sum ^ Bytes[I]) * 0x100000001b3ull;
+ for (; I + 8 <= Size; I += 8)
+ Sum = (Sum ^ *(const volatile u64 *)(Bytes + I)) * 0x100000001b3ull;
+ for (; I < Size; ++I)
+ Sum = (Sum ^ Bytes[I]) * 0x100000001b3ull;
+ return Sum;
+}
+
+static u64 ReadInstrumented(const volatile void *Ptr, uptr Size) {
+ switch (Size) {
+ case 1:
+ return *(const volatile u8 *)Ptr;
+ case 2:
+ return *(const volatile u16 *)Ptr;
+ case 4:
+ return *(const volatile u32 *)Ptr;
+ case 8:
+ return *(const volatile u64 *)Ptr;
+ default:
+ return ReadRange((const volatile u8 *)Ptr, Size);
+ }
+}
+
+static AccessInfo MakeAccessInfo(const volatile void *Ptr, uptr Size, int Type,
+ uptr PC, uptr BP) {
+ AccessInfo AI;
+ AI.ptr = Ptr;
+ AI.size = Size;
+ AI.access_type = Type;
+ AI.tid = (u32)GetTid();
+ AI.pc = PC;
+ AI.bp = BP;
+ return AI;
+}
+
+NOINLINE static void FoundWatchpoint(const volatile void *Ptr, uptr Size,
+ int Type, uptr PC, uptr, u64 *WP,
+ u64 Encoded) {
+ ScopedDisable Disable;
+ GetHostWatchpoints().TryConsume(WP, Encoded, PC,
+ (Type & CSAN_ACCESS_WRITE) != 0, Size);
+}
+
+NOINLINE static void SetupWatchpoint(const volatile void *Ptr, uptr Size,
+ int Type, uptr PC, uptr BP) {
+ ScopedDisable Disable;
+ ResetSkip();
+
+ if ((uptr)Ptr < GetPageSizeCached())
+ return;
+
+ u64 *WP = GetHostWatchpoints().Insert((uptr)Ptr, Size,
+ (Type & CSAN_ACCESS_WRITE) != 0,
+ HostWatchpointSlot((uptr)Ptr));
+ if (!WP)
+ return;
+
+ const u64 Old = ReadInstrumented(Ptr, Size);
+ DelayAccess(Type);
+ const u64 New = ReadInstrumented(Ptr, Size);
+
+ ValueChange VC = Old != New ? kValueChangeTrue : kValueChangeMaybe;
+
+ const AccessInfo AI = MakeAccessInfo(Ptr, Size, Type, PC, BP);
+ void *Peer;
+ int PeerAccess;
+ u32 PeerSize;
+ if (!GetHostWatchpoints().Consume(WP, Peer, PeerAccess, PeerSize)) {
+ ReportKnownOrigin(AI, VC, (uptr)Peer, PeerAccess, PeerSize, Old, New);
+ } else if (VC == kValueChangeTrue) {
+ ReportUnknownOrigin(AI, Old, New);
+ }
+
+ GetHostWatchpoints().Remove(WP);
+}
+
+ALWAYS_INLINE static void CheckAccess(const volatile void *Ptr, uptr Size,
+ int Type, uptr PC, uptr BP) {
+ if (UNLIKELY(!Ptr || !Size))
+ return;
+ if (Size > kHostMaxAccessSize)
+ Size = kHostMaxAccessSize;
+ if (UNLIKELY(uptr(Ptr) & ~HostWatchpointTable::AddressMask))
+ return;
+ Initialize();
+ if (UNLIKELY(DisableCount))
+ return;
+
+ u64 Encoded;
+ u64 *WP =
+ GetHostWatchpoints().Find((uptr)Ptr, Size, !(Type & CSAN_ACCESS_WRITE),
+ HostWatchpointSlot((uptr)Ptr), Encoded);
+ if (UNLIKELY(WP != nullptr))
+ FoundWatchpoint(Ptr, Size, Type, PC, BP, WP, Encoded);
+ else if (UNLIKELY(ShouldWatch(Type)))
+ SetupWatchpoint(Ptr, Size, Type, PC, BP);
+}
+
+static StaticSpinMutex InitMutex;
+static atomic_uint8_t Initialized;
+
+void Initialize() {
+ if (LIKELY(atomic_load(&Initialized, memory_order_acquire)))
+ return;
+ SpinMutexLock L(&InitMutex);
+ if (atomic_load(&Initialized, memory_order_relaxed))
+ return;
+ SanitizerToolName = "ConcurrencySanitizer";
+ CacheBinaryName();
+ InitializeFlags();
+ atomic_store(&Initialized, 1, memory_order_release);
+ Symbolizer::LateInitialize();
+}
+
+} // namespace __csan
+
+SANITIZER_INTERFACE_WEAK_DEF(const char *, __csan_default_options, void) {
+ return "";
+}
+
+INTERFACE u64 __csan_get_num_data_races() {
+ return atomic_load(&__csan::NumDataRaces, memory_order_relaxed);
+}
+
+INTERFACE void __csan_init() { __csan::Initialize(); }
+
+INTERFACE void __csan_func_entry(void *) {}
+INTERFACE void __csan_func_exit() {}
+INTERFACE void __csan_ignore_thread_begin() { ++__csan::DisableCount; }
+INTERFACE void __csan_ignore_thread_end() {
+ if (__csan::DisableCount)
+ --__csan::DisableCount;
+}
+
+static int AccessFlags(int Flags, bool IsWrite) {
+ return Flags | (IsWrite ? CSAN_ACCESS_WRITE : 0);
+}
+
+#define CSAN_PROBE(name, N, IsWrite) \
+ INTERFACE void name(void *Addr, int Flags) { \
+ GET_CALLER_PC_BP; \
+ __csan::CheckAccess(Addr, N, AccessFlags(Flags, IsWrite), pc, bp); \
+ }
+
+#define CSAN_ACCESS(N) \
+ CSAN_PROBE(__csan_read##N, N, false) \
+ CSAN_PROBE(__csan_unaligned_read##N, N, false) \
+ CSAN_PROBE(__csan_volatile_read##N, N, false) \
+ CSAN_PROBE(__csan_unaligned_volatile_read##N, N, false) \
+ CSAN_PROBE(__csan_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_write##N, N, true) \
+ CSAN_PROBE(__csan_volatile_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_volatile_write##N, N, true) \
+ CSAN_PROBE(__csan_read_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_read_write##N, N, true)
+
+CSAN_ACCESS(1)
+CSAN_ACCESS(2)
+CSAN_ACCESS(4)
+CSAN_ACCESS(8)
+CSAN_ACCESS(16)
+
+INTERFACE void __csan_read_range(void *Addr, uptr Size, int Flags) {
+ GET_CALLER_PC_BP;
+ __csan::CheckAccess(Addr, Size, AccessFlags(Flags, false), pc, bp);
+}
+
+INTERFACE void __csan_write_range(void *Addr, uptr Size, int Flags) {
+ GET_CALLER_PC_BP;
+ __csan::CheckAccess(Addr, Size, AccessFlags(Flags, true), pc, bp);
+}
+
+// TODO: Handle thread reordering checks like KCSAN.
+INTERFACE void __csan_atomic_thread_fence(int) {}
+INTERFACE void __csan_atomic_signal_fence(int) {}
diff --git a/compiler-rt/lib/csan/csan.h b/compiler-rt/lib/csan/csan.h
new file mode 100644
index 00000000000000..47d1eb046429ff
--- /dev/null
+++ b/compiler-rt/lib/csan/csan.h
@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Internal declarations for the host ConcurrencySanitizer runtime.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_H
+#define CSAN_H
+
+#include "csan_defs.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+
+namespace __csan {
+
+using __sanitizer::s32;
+using __sanitizer::u32;
+using __sanitizer::u64;
+using __sanitizer::uptr;
+
+struct Flags {
+#define CSAN_FLAG(Type, Name, DefaultValue, Description) Type Name;
+#include "csan_flags.inc"
+#undef CSAN_FLAG
+ void SetDefaults();
+};
+
+extern Flags flags_data;
+inline Flags *flags() { return &flags_data; }
+
+void InitializeFlags();
+
+enum ValueChange { kValueChangeMaybe, kValueChangeFalse, kValueChangeTrue };
+
+static constexpr u32 kMaxStackFrames = 64;
+
+struct AccessInfo {
+ const volatile void *ptr;
+ uptr size;
+ int access_type;
+ u32 tid;
+ uptr pc;
+ uptr bp;
+};
+
+void RecordDataRace();
+void ReportKnownOrigin(const AccessInfo &AI, ValueChange VC, uptr PeerPC,
+ int PeerAccess, uptr PeerSize, u64 Old, u64 New);
+void ReportUnknownOrigin(const AccessInfo &AI, u64 Old, u64 New);
+
+} // namespace __csan
+
+extern "C" {
+SANITIZER_INTERFACE_ATTRIBUTE SANITIZER_WEAK_ATTRIBUTE const char *
+__csan_default_options();
+}
+
+#endif // CSAN_H
diff --git a/compiler-rt/lib/csan/csan_defs.h b/compiler-rt/lib/csan/csan_defs.h
new file mode 100644
index 00000000000000..e442b2b7e798c6
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_defs.h
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Constants shared by the host and device ConcurrencySanitizer runtimes.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_DEFS_H
+#define CSAN_DEFS_H
+
+enum {
+ CSAN_ACCESS_ATOMIC = 1u << 0,
+ CSAN_ACCESS_COMPOUND = 1u << 1,
+ CSAN_ACCESS_WRITE = 1u << 2
+};
+
+enum {
+ CSAN_RACE_DATA = 0,
+ CSAN_RACE_UNKNOWN_ORIGIN = 1,
+ CSAN_RACE_INTRA_WAVE = 2
+};
+
+// The size of the device watchpoint table used by the host and device runtime.
+static constexpr unsigned long CSAN_WATCHPOINT_TABLE_BYTES =
+ 2ul * 1024ul * 1024ul;
+static constexpr unsigned long CSAN_WATCHPOINT_TABLE_ENTRIES =
+ CSAN_WATCHPOINT_TABLE_BYTES / sizeof(unsigned long long);
+
+#endif // CSAN_DEFS_H
diff --git a/compiler-rt/lib/csan/csan_flags.inc b/compiler-rt/lib/csan/csan_flags.inc
new file mode 100644
index 00000000000000..692861f37d87f4
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_flags.inc
@@ -0,0 +1,23 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Host ConcurrencySanitizer runtime flags.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_FLAG
+#error "Define CSAN_FLAG prior to including this file!"
+#endif
+
+// CSAN_FLAG(Type, Name, DefaultValue, Description)
+
+CSAN_FLAG(int, udelay, 80, "Microseconds to stall after arming a watchpoint.")
+CSAN_FLAG(int, skip_watch, 4000,
+ "Skip this many accesses (per thread) before arming a watchpoint.")
+CSAN_FLAG(bool, halt_on_error, false, "Die after the first reported race.")
diff --git a/compiler-rt/lib/csan/csan_gpu.cpp b/compiler-rt/lib/csan/csan_gpu.cpp
new file mode 100644
index 00000000000000..729113f0c7950e
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_gpu.cpp
@@ -0,0 +1,456 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Watchpoint-based data race detector for GPU targets.
+///
+//===----------------------------------------------------------------------===//
+
+#include <gpuintrin.h>
+
+#include "csan_offload_packet.h"
+#include "csan_watch.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+#include "shared/rpc.h"
+
+using namespace __sanitizer;
+
+#define INTERFACE extern "C" SANITIZER_INTERFACE_ATTRIBUTE
+
+extern "C" {
+// Externally initialized by the sanitizer, keeps one table per active device.
+[[gnu::visibility("protected")]] u64 *__csan_watchpoint_table = nullptr;
+}
+
+static constexpr u64 SAMPLE_DELAY_MIN_NS = 1000;
+static constexpr u64 SAMPLE_DELAY_MAX_NS = 10000;
+static constexpr u32 WP_CHANCE = 8;
+static_assert((CSAN_WATCHPOINT_TABLE_ENTRIES &
+ (CSAN_WATCHPOINT_TABLE_ENTRIES - 1)) == 0,
+ "watchpoint table size must be a power of two");
+static_assert(WP_CHANCE >= 2 && (WP_CHANCE & (WP_CHANCE - 1)) == 0,
+ "WP_CHANCE must be a power of two");
+
+// The GPU case does
+static constexpr u32 GPU_MAX_ACCESS_SIZE = 16;
+static constexpr u32 GPU_WATCHPOINT_ENTRIES = CSAN_WATCHPOINT_TABLE_ENTRIES;
+static constexpr u32 GPU_CHECK_ADJACENT_SLOTS = 0;
+static_assert(GPU_WATCHPOINT_ENTRIES * sizeof(u64) == 2 * 1024 * 1024,
+ "GPU watchpoint table must be 2 MiB");
+using GpuWatchpointTable =
+ __csan::WatchpointTable<GPU_MAX_ACCESS_SIZE, GPU_CHECK_ADJACENT_SLOTS>;
+
+static GpuWatchpointTable get_watchpoints() {
+ return GpuWatchpointTable(__csan_watchpoint_table);
+}
+
+// LDS addresses share the global watchpoint table. The original address is
+// combined with the block's linear ID to create a unique global address.
+static constexpr u32 LDS_OFFSET_BITS = 20;
+static constexpr u64 LDS_OFFSET_MASK = (1ull << LDS_OFFSET_BITS) - 1;
+static constexpr u64 LDS_FLAG = 1ull << (GpuWatchpointTable::AddressBits - 1);
+static constexpr u64 GLOBAL_ADDRESS_MASK = LDS_FLAG - 1;
+static constexpr u64 LDS_MAX_BLOCKS =
+ 1ull << (GpuWatchpointTable::AddressBits - 1 - LDS_OFFSET_BITS);
+static_assert((LDS_FLAG | ((LDS_MAX_BLOCKS - 1) << LDS_OFFSET_BITS) |
+ LDS_OFFSET_MASK) == GpuWatchpointTable::AddressMask,
+ "LDS key fields must fill the address bits");
+
+[[gnu::visibility("protected"),
+ gnu::weak]] rpc::Client client asm("__llvm_rpc_client");
+
+static u64 __csan_num_data_races = 0;
+
+INTERFACE u64 __csan_get_num_data_races() {
+ return __atomic_load_n(&__csan_num_data_races, __ATOMIC_RELAXED);
+}
+
+// Shallow deduplication check to save the host thread work. Keyed on both the
+// PC and the race kind so each distinct kind of race at a PC is reported once.
+static bool should_report(void *pc, unsigned kind) {
+ static u64 seen[64] = {};
+ const u64 token = (reinterpret_cast<uptr>(pc) >> 4) ^
+ (static_cast<u64>(kind) * 0x9E3779B97F4A7C15ull);
+ u64 idx = (token * 0x9E3779B97F4A7C15ull) >> 58;
+ u64 last = __scoped_atomic_exchange_n(&seen[idx], token, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+ return last != token;
+}
+
+[[gnu::cold, gnu::noinline]] static void
+report(unsigned kind, uptr addr, u32 size, int access_type, uptr pc,
+ void *peer = nullptr, int peer_access = 0, u32 peer_size = 0,
+ u8 peer_lane = 0) {
+ pc = pc ? pc : GET_CALLER_PC();
+ if (!should_report(reinterpret_cast<void *>(pc), kind))
+ return;
+
+ __csan_gpu_race rep = {};
+ rep.pc = pc;
+ rep.peer_pc = reinterpret_cast<uptr>(peer);
+ rep.addr = addr;
+ rep.size = size;
+ rep.access_type = static_cast<unsigned>(access_type);
+ rep.kind = kind;
+ rep.block[0] = __gpu_block_id(__GPU_X_DIM);
+ rep.block[1] = __gpu_block_id(__GPU_Y_DIM);
+ rep.block[2] = __gpu_block_id(__GPU_Z_DIM);
+ rep.thread[0] = __gpu_thread_id(__GPU_X_DIM);
+ rep.thread[1] = __gpu_thread_id(__GPU_Y_DIM);
+ rep.thread[2] = __gpu_thread_id(__GPU_Z_DIM);
+ rep.lane = __gpu_lane_id();
+ rep.peer_lane = peer_lane;
+ rep.peer_access_type = static_cast<u8>(peer_access);
+ rep.peer_size = static_cast<u8>(peer_size);
+
+ rpc::Client::Port Port = client.open<SANITIZER_OFFLOAD_CSAN>();
+ Port.send([&](rpc::Buffer *buf, u32) {
+ __builtin_memcpy(buf->data, &rep, sizeof(rep));
+ });
+ static_assert(sizeof(__csan_gpu_race) <= sizeof(rpc::Buffer),
+ "Report must fit in a single packet");
+
+ __scoped_atomic_fetch_add(&__csan_num_data_races, 1, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+}
+
+#if defined(__AMDGPU__)
+// AMDGPU does not have a single set frequency. Different architectures and
+// cards can have different values. A frequency of 100MHz is most common so we
+// use it, if it is wrong it just means we sleep longer than expected.
+static constexpr u64 CLOCK_FREQ_HZ = 100000000UL;
+#else
+static constexpr u64 CLOCK_FREQ_HZ = 1000000000UL;
+#endif
+static constexpr u64 TICKS_PER_SEC = 1000000000UL;
+
+// FIXME: Avoids emitting an unresolved reference to the OCLC ABI version.
+static u32 num_blocks(int dim) {
+#ifdef __AMDGPU__
+ return ((const u32 __gpu_constant *)__builtin_amdgcn_implicitarg_ptr())[dim];
+#else
+ return __gpu_num_blocks(dim);
+#endif
+}
+
+static u64 lds_block_index() {
+ return (u64)__gpu_block_id(__GPU_X_DIM) +
+ (u64)num_blocks(__GPU_X_DIM) *
+ ((u64)__gpu_block_id(__GPU_Y_DIM) +
+ (u64)num_blocks(__GPU_Y_DIM) * (u64)__gpu_block_id(__GPU_Z_DIM));
+}
+
+static bool lds_fits() {
+ u64 nxy, nxyz;
+ if (__builtin_mul_overflow((u64)num_blocks(__GPU_X_DIM),
+ (u64)num_blocks(__GPU_Y_DIM), &nxy) ||
+ __builtin_mul_overflow(nxy, (u64)num_blocks(__GPU_Z_DIM), &nxyz))
+ return false;
+ return nxyz <= LDS_MAX_BLOCKS;
+}
+
+// Stateless PRNG hashes cycle counter and global thread ID using SplitMix64.
+static u64 rng() {
+ u64 z = __builtin_readcyclecounter();
+ z ^= (u64(__gpu_block_id(__GPU_X_DIM)) << 32 | __gpu_thread_id(__GPU_X_DIM)) *
+ 0xD1B54A32D192ED03ull;
+ z += 0x9E3779B97F4A7C15ull;
+ z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull;
+ z = (z ^ (z >> 27)) * 0x94D049BB133111EBull;
+ return z ^ (z >> 31);
+}
+
+namespace {
+template <typename> struct is_ptr_local {
+ static constexpr bool value = false;
+};
+template <typename T> struct is_ptr_local<T __gpu_local *> {
+ static constexpr bool value = true;
+};
+} // namespace
+
+template <typename PtrTy> static uptr report_address(PtrTy addr) {
+ if constexpr (is_ptr_local<PtrTy>::value)
+ return reinterpret_cast<uptr>((const volatile void *)addr);
+ return reinterpret_cast<uptr>(addr);
+}
+
+template <typename PtrTy> static bool uses_table() {
+ if constexpr (is_ptr_local<PtrTy>::value)
+ return lds_fits();
+ return true;
+}
+
+template <typename PtrTy> static u64 watch_key(uptr addr) {
+ if constexpr (is_ptr_local<PtrTy>::value)
+ return LDS_FLAG | (lds_block_index() << LDS_OFFSET_BITS) |
+ (addr & LDS_OFFSET_MASK);
+ return addr & GLOBAL_ADDRESS_MASK;
+}
+
+static u32 watchpoint_slot(u64 key) {
+ key ^= key >> LDS_OFFSET_BITS;
+ return (key / GPU_MAX_ACCESS_SIZE) & (GPU_WATCHPOINT_ENTRIES - 1);
+}
+
+static bool should_watch(u64 lane_mask, u32 access_type) {
+ // If every access is atomic we cannot have a race.
+ if (!__gpu_ballot(lane_mask, !(access_type & CSAN_ACCESS_ATOMIC)))
+ return false;
+
+ constexpr unsigned N = __builtin_ctzg(WP_CHANCE);
+ return __gpu_read_first_lane_u32(lane_mask, (rng() >> (64 - N)) == 0);
+}
+
+template <typename PtrTy>
+static u64 *find_watchpoint(uptr addr, u32 size, bool expect_write,
+ u64 &encoded) {
+ const u64 key = watch_key<PtrTy>(addr);
+ return get_watchpoints().Find(key, size, expect_write, watchpoint_slot(key),
+ encoded);
+}
+
+// FNV-1a digest of a byte range so wide accesses can reuse the value comparison
+// semantics.
+template <typename BytePtr, typename WordPtr>
+static u64 read_range(BytePtr bytes, WordPtr, u32 size) {
+ u64 sum = 0xcbf29ce484222325ull;
+ u32 i = 0;
+
+ for (; i < size && ((reinterpret_cast<uptr>(bytes) + i) & 7u); ++i)
+ sum = (sum ^ bytes[i]) * 0x100000001b3ull;
+ for (; i + 8 <= size; i += 8)
+ sum = (sum ^ *reinterpret_cast<WordPtr>(bytes + i)) * 0x100000001b3ull;
+ for (; i < size; ++i)
+ sum = (sum ^ bytes[i]) * 0x100000001b3ull;
+ return sum;
+}
+
+// Snapshot the watched location for value-change detection. Larger sizes get
+// converted into a single checksum.
+static u64 read_instrumented_memory(const volatile __gpu_global void *ptr,
+ u32 size) {
+ const uptr addr =
+ reinterpret_cast<uptr>(const_cast<const __gpu_global void *>(ptr));
+ if ((addr & (size - 1)) == 0) {
+ switch (size) {
+ case 1:
+ return *(const volatile __gpu_global u8 *)ptr;
+ case 2:
+ return *(const volatile __gpu_global u16 *)ptr;
+ case 4:
+ return *(const volatile __gpu_global u32 *)ptr;
+ case 8:
+ return *(const volatile __gpu_global u64 *)ptr;
+ }
+ }
+ return read_range((const volatile __gpu_global u8 *)ptr,
+ (const volatile __gpu_global u64 *)ptr, size);
+}
+
+static u64 read_instrumented_memory(const volatile __gpu_local void *ptr,
+ u32 size) {
+ const uptr addr =
+ reinterpret_cast<uptr>(const_cast<const __gpu_local void *>(ptr));
+ if ((addr & (size - 1)) == 0) {
+ switch (size) {
+ case 1:
+ return *(const volatile __gpu_local u8 *)ptr;
+ case 2:
+ return *(const volatile __gpu_local u16 *)ptr;
+ case 4:
+ return *(const volatile __gpu_local u32 *)ptr;
+ case 8:
+ return *(const volatile __gpu_local u64 *)ptr;
+ }
+ }
+ return read_range((const volatile __gpu_local u8 *)ptr,
+ (const volatile __gpu_local u64 *)ptr, size);
+}
+
+static bool intra_wave_race(u64 lane_mask, uptr addr, int access_type,
+ u8 &peer_lane) {
+ const bool is_write = (access_type & CSAN_ACCESS_WRITE) != 0;
+ const bool is_atomic = (access_type & CSAN_ACCESS_ATOMIC) != 0;
+ const u64 writers = __gpu_ballot(lane_mask, is_write);
+ const u64 nonatomic = __gpu_ballot(lane_mask, !is_atomic);
+ if (!writers || !nonatomic)
+ return false;
+
+ const u64 same_addr = __gpu_match_any_u64(lane_mask, addr);
+ const bool is_race = __builtin_popcountg(same_addr) >= 2 &&
+ (same_addr & writers) && (same_addr & nonatomic);
+ if (!is_race || !__gpu_is_first_in_lane(same_addr))
+ return false;
+ peer_lane = static_cast<u8>(63u - __builtin_clzg(same_addr));
+ return true;
+}
+
+static void delay_ns(u64 nsecs) {
+ const u64 tick_rate = TICKS_PER_SEC / CLOCK_FREQ_HZ;
+ const u64 start = __builtin_readsteadycounter();
+ const u64 end = start + (nsecs + tick_rate - 1) / tick_rate;
+#if defined(__AMDGPU__)
+ __builtin_amdgcn_s_sleep(2);
+ while (__builtin_readsteadycounter() < end)
+ __builtin_amdgcn_s_sleep(15);
+#else
+ while (__builtin_readsteadycounter() < end)
+ __gpu_thread_suspend();
+#endif
+}
+
+static void sample_delay(u64 lane_mask) {
+ u64 nsecs = SAMPLE_DELAY_MIN_NS;
+ if (__gpu_is_first_in_lane(lane_mask))
+ nsecs += (rng() >> 32) % (SAMPLE_DELAY_MAX_NS - SAMPLE_DELAY_MIN_NS);
+ delay_ns(__gpu_read_first_lane_u64(lane_mask, nsecs));
+}
+
+[[gnu::cold, gnu::noinline]] static void
+found_watchpoint(u64 *wp, u64 encoded, uptr pc, bool is_write, u32 size) {
+ pc = pc ? pc : GET_CALLER_PC();
+ get_watchpoints().TryConsume(wp, encoded, pc, is_write, size);
+}
+
+// The slow path, sets a watchpoint in the table and waits to see if any other
+// thread tripped it. Returns a CSAN_RACE_* kind, or -1 if none.
+template <typename PtrTy>
+static int watch(u64 lane_mask, const PtrTy addr, u32 size, int access_type,
+ uptr pc, uptr &report_pc, void *&peer, int &peer_access,
+ u32 &peer_size) {
+ report_pc = pc ? pc : GET_CALLER_PC();
+ const bool is_write = (access_type & CSAN_ACCESS_WRITE) != 0;
+ const uptr iaddr = reinterpret_cast<uptr>(addr);
+
+ const u32 wp_size = GpuWatchpointTable::EncodeSize(size);
+ const bool armable = uses_table<PtrTy>() &&
+ !(access_type & CSAN_ACCESS_ATOMIC) &&
+ (is_ptr_local<PtrTy>::value || iaddr != 0);
+ const u64 key = watch_key<PtrTy>(iaddr);
+ u64 *wp = armable ? get_watchpoints().Insert(key, wp_size, is_write,
+ watchpoint_slot(key))
+ : nullptr;
+
+ const u64 old = read_instrumented_memory(addr, size);
+ sample_delay(lane_mask);
+ const u64 now = read_instrumented_memory(addr, size);
+
+ peer = nullptr;
+ peer_access = 0;
+ peer_size = 0;
+ int kind = -1;
+ if (wp && !get_watchpoints().Consume(wp, peer, peer_access, peer_size))
+ kind = CSAN_RACE_DATA;
+ else if (old != now)
+ kind = CSAN_RACE_UNKNOWN_ORIGIN;
+
+ if (wp)
+ get_watchpoints().Remove(wp);
+ return kind;
+}
+
+template <typename PtrTy>
+static void check_access_impl(u64 lane_mask, const PtrTy addr, u32 size,
+ int access_type, uptr pc) {
+ pc = pc ? pc : GET_CALLER_PC();
+ if (uses_table<PtrTy>()) {
+ const bool is_write = (access_type & CSAN_ACCESS_WRITE) != 0;
+ u64 encoded;
+ u64 *wp = find_watchpoint<PtrTy>((u64)addr, size, !is_write, encoded);
+ if (wp)
+ found_watchpoint(wp, encoded, pc, is_write, size);
+ }
+
+ if (!should_watch(lane_mask, access_type))
+ return;
+
+ const uptr iaddr = reinterpret_cast<uptr>(addr);
+ u8 peer_lane;
+ if (intra_wave_race(lane_mask, iaddr, access_type, peer_lane))
+ report(CSAN_RACE_INTRA_WAVE, report_address(addr), size, access_type, pc,
+ nullptr, 0, 0, peer_lane);
+
+ if (access_type & CSAN_ACCESS_ATOMIC)
+ return;
+
+ uptr report_pc;
+ void *peer;
+ int peer_access;
+ u32 peer_size;
+ int kind = watch(lane_mask, addr, size, access_type, pc, report_pc, peer,
+ peer_access, peer_size);
+ if (kind >= 0)
+ report(static_cast<unsigned>(kind), report_address(addr), size, access_type,
+ report_pc, peer, peer_access, peer_size);
+}
+
+static void check_access(const volatile void *addr, uptr size, int access_type,
+ uptr pc) {
+ pc = pc ? pc : GET_CALLER_PC();
+ if (__gpu_is_ptr_private(const_cast<void *>(addr)) || !size)
+ return;
+ if (size > ~u32(0))
+ size = ~u32(0);
+
+ if (__gpu_is_ptr_local(const_cast<void *>(addr)))
+ return check_access_impl(__gpu_lane_mask(),
+ (const volatile __gpu_local void *)addr, size,
+ access_type, pc);
+ check_access_impl(__gpu_lane_mask(), (const volatile __gpu_global void *)addr,
+ size, access_type, pc);
+}
+
+//===----------------------------------------------------------------------===//
+// Public ABI (emitted by the ConcurrencySanitizer pass)
+//===----------------------------------------------------------------------===//
+
+// Using `sanitize_concurrency_no_checking_at_run_time` ignored on the device,
+INTERFACE void __csan_init() {}
+INTERFACE void __csan_func_entry(void *) {}
+INTERFACE void __csan_func_exit() {}
+INTERFACE void __csan_ignore_thread_begin() {}
+INTERFACE void __csan_ignore_thread_end() {}
+
+static int access_flags(int flags, bool is_write) {
+ return flags | (is_write ? CSAN_ACCESS_WRITE : 0);
+}
+
+#define CSAN_PROBE(name, N, is_write) \
+ INTERFACE void name(void *addr, int flags) { \
+ check_access(addr, N, access_flags(flags, is_write), GET_CALLER_PC()); \
+ }
+
+#define CSAN_ACCESS(N) \
+ CSAN_PROBE(__csan_read##N, N, false) \
+ CSAN_PROBE(__csan_unaligned_read##N, N, false) \
+ CSAN_PROBE(__csan_volatile_read##N, N, false) \
+ CSAN_PROBE(__csan_unaligned_volatile_read##N, N, false) \
+ CSAN_PROBE(__csan_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_write##N, N, true) \
+ CSAN_PROBE(__csan_volatile_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_volatile_write##N, N, true) \
+ CSAN_PROBE(__csan_read_write##N, N, true) \
+ CSAN_PROBE(__csan_unaligned_read_write##N, N, true)
+
+CSAN_ACCESS(1)
+CSAN_ACCESS(2)
+CSAN_ACCESS(4)
+CSAN_ACCESS(8)
+CSAN_ACCESS(16)
+
+INTERFACE void __csan_read_range(void *addr, uptr size, int flags) {
+ check_access(addr, size, access_flags(flags, false), GET_CALLER_PC());
+}
+
+INTERFACE void __csan_write_range(void *addr, uptr size, int flags) {
+ check_access(addr, size, access_flags(flags, true), GET_CALLER_PC());
+}
+
+INTERFACE void __csan_atomic_thread_fence(int) {}
+INTERFACE void __csan_atomic_signal_fence(int) {}
diff --git a/compiler-rt/lib/csan/csan_offload_packet.h b/compiler-rt/lib/csan/csan_offload_packet.h
new file mode 100644
index 00000000000000..827eaf06ebfada
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_offload_packet.h
@@ -0,0 +1,39 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// RPC packet shared by the device and host ConcurrencySanitizer runtimes.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_OFFLOAD_PACKET_H
+#define CSAN_OFFLOAD_PACKET_H
+
+#include "csan_defs.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+#include "sanitizer_common/sanitizer_offload_opcodes.h"
+
+struct __csan_gpu_race {
+ __sanitizer::u64 pc;
+ __sanitizer::u64 peer_pc;
+ __sanitizer::u64 addr;
+ __sanitizer::u32 size;
+ __sanitizer::u32 access_type;
+ __sanitizer::u32 kind;
+ __sanitizer::u32 block[3];
+ __sanitizer::u16 thread[3];
+ __sanitizer::u8 lane;
+ __sanitizer::u8 peer_lane;
+ __sanitizer::u8 peer_access_type;
+ __sanitizer::u8 peer_size;
+};
+
+static_assert(sizeof(__csan_gpu_race) == 64,
+ "Offload CSan report must fit one RPC packet");
+
+#endif // CSAN_OFFLOAD_PACKET_H
diff --git a/compiler-rt/lib/csan/csan_report.cpp b/compiler-rt/lib/csan/csan_report.cpp
new file mode 100644
index 00000000000000..3a641c307f3871
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_report.cpp
@@ -0,0 +1,279 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Host ConcurrencySanitizer report generation.
+///
+//===----------------------------------------------------------------------===//
+
+#include "csan.h"
+
+#include "sanitizer_common/sanitizer_common.h"
+#include "sanitizer_common/sanitizer_flags.h"
+#include "sanitizer_common/sanitizer_libc.h"
+#include "sanitizer_common/sanitizer_mutex.h"
+#include "sanitizer_common/sanitizer_report_decorator.h"
+#include "sanitizer_common/sanitizer_stacktrace.h"
+#include "sanitizer_common/sanitizer_symbolizer.h"
+
+using namespace __sanitizer;
+
+namespace __sanitizer {
+void BufferedStackTrace::UnwindImpl(uptr pc, uptr bp, void *context,
+ bool request_fast, u32 max_depth) {
+ uptr top = 0;
+ uptr bottom = 0;
+ GetThreadStackTopAndBottom(false, &top, &bottom);
+ bool fast = StackTrace::WillUseFastUnwind(request_fast);
+ Unwind(max_depth, pc, bp, context, top, bottom, fast);
+}
+} // namespace __sanitizer
+
+namespace __csan {
+namespace {
+
+class Decorator : public SanitizerCommonDecorator {
+public:
+ const char *Access() { return Blue(); }
+ const char *Location() { return Green(); }
+};
+
+// A pair of racing program counter values to deduplicate and check.
+struct RacyPcs {
+ uptr pc[2];
+
+ bool operator==(const RacyPcs &Other) const {
+ if (pc[0] == Other.pc[0] && pc[1] == Other.pc[1])
+ return true;
+ return pc[0] == Other.pc[1] && pc[1] == Other.pc[0];
+ }
+};
+InternalMmapVectorNoCtor<RacyPcs> RacyPcsSeen;
+
+struct PeerInfo {
+ uptr pc;
+ uptr size;
+ int access_type;
+};
+
+bool HandleRacyPcs(uptr PC, uptr PeerPC) {
+ const RacyPcs Racy = {{PC, PeerPC}};
+ for (uptr I = 0; I < RacyPcsSeen.size(); ++I) {
+ if (Racy == RacyPcsSeen[I]) {
+ VReport(2, "%s: suppressing report as doubled\n", SanitizerToolName);
+ return true;
+ }
+ }
+ RacyPcsSeen.push_back(Racy);
+ return false;
+}
+
+const char *AccessKind(int Type) {
+ const bool Write = Type & CSAN_ACCESS_WRITE;
+ const bool Atomic = Type & CSAN_ACCESS_ATOMIC;
+ const bool Compound = Type & CSAN_ACCESS_COMPOUND;
+ if (Compound && Write)
+ return Atomic ? "read-write (atomic)" : "read-write";
+ if (Write)
+ return Atomic ? "write (atomic)" : "write";
+ return Atomic ? "read (atomic)" : "read";
+}
+
+const char *MemOpDesc(bool First, int Type) {
+ const bool Write = Type & CSAN_ACCESS_WRITE;
+ const bool Atomic = Type & CSAN_ACCESS_ATOMIC;
+ const bool Compound = Type & CSAN_ACCESS_COMPOUND;
+ if (Compound && Write)
+ return Atomic ? (First ? "Read-write (atomic)"
+ : "Previous read-write (atomic)")
+ : (First ? "Read-write" : "Previous read-write");
+ if (Write)
+ return Atomic ? (First ? "Write (atomic)" : "Previous write (atomic)")
+ : (First ? "Write" : "Previous write");
+ return Atomic ? (First ? "Read (atomic)" : "Previous read (atomic)")
+ : (First ? "Read" : "Previous read");
+}
+
+void CaptureStack(uptr PC, uptr BP, uptr *Out, u32 *N) {
+ UNINITIALIZED BufferedStackTrace Stack;
+ Stack.Unwind(PC, BP, nullptr, common_flags()->fast_unwind_on_fatal,
+ kMaxStackFrames);
+ *N = Stack.size > kMaxStackFrames ? kMaxStackFrames : Stack.size;
+ if (*N)
+ internal_memcpy(Out, Stack.trace, *N * sizeof(uptr));
+}
+
+SymbolizedStack *SymbolizeFrame(uptr PC) {
+ return Symbolizer::GetOrInit()->SymbolizePC(
+ StackTrace::GetPreviousInstructionPc(PC));
+}
+
+u32 SkipRuntimeFrames(const uptr *PCs, u32 N) {
+ for (u32 I = 0; I < N; ++I) {
+ SymbolizedStack *Frames = SymbolizeFrame(PCs[I]);
+ const bool User = Frames && SkipInternalFrames(Frames);
+ if (Frames)
+ Frames->ClearAll();
+ if (User)
+ return I;
+ }
+ return 0;
+}
+
+void CopyFuncName(uptr PC, InternalScopedString *Out) {
+ SymbolizedStack *Frames = SymbolizeFrame(PC);
+ const SymbolizedStack *User = Frames ? SkipInternalFrames(Frames) : nullptr;
+ if (!User)
+ User = Frames;
+ if (User && User->info.function && User->info.function[0])
+ Out->Append(User->info.function);
+ else
+ Out->AppendF("%p", (void *)PC);
+ if (Frames)
+ Frames->ClearAll();
+}
+
+void PrintFrames(const uptr *PCs, u32 N) {
+ if (!N) {
+ Printf(" [failed to restore the stack]\n\n");
+ return;
+ }
+ const u32 Skip = SkipRuntimeFrames(PCs, N);
+ StackTrace Trace(PCs + Skip, N - Skip);
+ Trace.Print();
+}
+
+void PrintLocation(uptr Addr) {
+ DataInfo Loc;
+ if (!Symbolizer::GetOrInit()->SymbolizeData(Addr, &Loc) || !Loc.name ||
+ !Loc.name[0] || Loc.name[0] == '?') {
+ Loc.Clear();
+ return;
+ }
+ Decorator D;
+ Printf("%s", D.Location());
+ if (Loc.size)
+ Printf(" Location is global '%s' of size %zu at %p\n", Loc.name, Loc.size,
+ (void *)Addr);
+ else
+ Printf(" Location is global '%s' at %p\n", Loc.name, (void *)Addr);
+ Printf("%s", D.Default());
+ Loc.Clear();
+}
+
+void PrintHexValue(u64 V, uptr Size) {
+ switch (Size) {
+ case 1:
+ Printf("0x%02llx", (unsigned long long)V);
+ break;
+ case 2:
+ Printf("0x%04llx", (unsigned long long)V);
+ break;
+ case 4:
+ Printf("0x%08llx", (unsigned long long)V);
+ break;
+ default:
+ Printf("0x%016llx", (unsigned long long)V);
+ break;
+ }
+}
+
+void PrintValueChange(uptr Size, u64 Old, u64 New) {
+ if (Size == 0 || Size > 8 || Old == New)
+ return;
+ Printf(" value changed: ");
+ PrintHexValue(Old, Size);
+ Printf(" -> ");
+ PrintHexValue(New, Size);
+ Printf("\n");
+}
+
+void PrintReport(const AccessInfo &AI, const PeerInfo *Other, u64 Old,
+ u64 New) {
+ if (HandleRacyPcs(AI.pc, Other ? Other->pc : 0))
+ return;
+
+ UNINITIALIZED uptr ThisStack[kMaxStackFrames];
+ u32 ThisN = 0;
+ CaptureStack(AI.pc, AI.bp, ThisStack, &ThisN);
+
+ const u32 ThisSkip = SkipRuntimeFrames(ThisStack, ThisN);
+ const uptr ThisFrame = ThisN ? ThisStack[ThisSkip] : AI.pc;
+ const uptr OtherFrame = Other ? Other->pc : 0;
+
+ RecordDataRace();
+
+ Decorator D;
+ InternalScopedString ThisFn;
+ CopyFuncName(ThisFrame, &ThisFn);
+
+ Printf("==================\n");
+ Printf("%s", D.Warning());
+ if (Other) {
+ InternalScopedString OtherFn;
+ CopyFuncName(OtherFrame, &OtherFn);
+ const int Cmp = internal_strcmp(OtherFn.data(), ThisFn.data());
+ Printf("WARNING: ConcurrencySanitizer: data race in %s / %s\n",
+ Cmp < 0 ? OtherFn.data() : ThisFn.data(),
+ Cmp < 0 ? ThisFn.data() : OtherFn.data());
+ } else {
+ Printf("WARNING: ConcurrencySanitizer: data race of unknown origin in %s\n",
+ ThisFn.data());
+ }
+ Printf("%s", D.Default());
+
+ Printf("%s", D.Access());
+ if (Other) {
+ Printf(" %s of size %zu at %p by thread %u:\n",
+ MemOpDesc(true, AI.access_type), AI.size, AI.ptr, AI.tid);
+ Printf("%s", D.Default());
+ PrintFrames(ThisStack, ThisN);
+
+ Printf("%s", D.Access());
+ Printf(" %s of size %zu at %p:\n", MemOpDesc(false, Other->access_type),
+ Other->size, AI.ptr);
+ Printf("%s", D.Default());
+ PrintFrames(&Other->pc, 1);
+ } else {
+ Printf(
+ " race at unknown origin, with %s of size %zu at %p by thread %u:\n",
+ AccessKind(AI.access_type), AI.size, AI.ptr, AI.tid);
+ Printf("%s", D.Default());
+ PrintFrames(ThisStack, ThisN);
+ }
+
+ PrintValueChange(AI.size, Old, New);
+ PrintLocation((uptr)AI.ptr);
+
+ if (ThisN) {
+ StackTrace Summary(ThisStack + ThisSkip, ThisN - ThisSkip);
+ ReportErrorSummary(Other ? "data race" : "data race of unknown origin",
+ &Summary);
+ }
+ Printf("==================\n");
+ if (flags()->halt_on_error)
+ Die();
+}
+
+} // namespace
+
+void ReportKnownOrigin(const AccessInfo &AI, ValueChange VC, uptr PeerPC,
+ int PeerAccess, uptr PeerSize, u64 Old, u64 New) {
+ ScopedErrorReportLock L;
+ if (VC == kValueChangeFalse)
+ return;
+ const PeerInfo Peer = {PeerPC, PeerSize, PeerAccess};
+ PrintReport(AI, &Peer, Old, New);
+}
+
+void ReportUnknownOrigin(const AccessInfo &AI, u64 Old, u64 New) {
+ ScopedErrorReportLock L;
+ PrintReport(AI, nullptr, Old, New);
+}
+
+} // namespace __csan
diff --git a/compiler-rt/lib/csan/csan_watch.h b/compiler-rt/lib/csan/csan_watch.h
new file mode 100644
index 00000000000000..6b7fff0593b12a
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_watch.h
@@ -0,0 +1,163 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Configurable packed watchpoint table shared by the host and GPU runtimes.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_WATCH_H
+#define CSAN_WATCH_H
+
+#include "csan_defs.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+
+#if !defined(__has_builtin) || !__has_builtin(__scoped_atomic_load_n)
+#define CSAN_DEFINED_SCOPED_ATOMICS
+#define __scoped_atomic_load_n(P, Order, Scope) __atomic_load_n(P, Order)
+#define __scoped_atomic_store_n(P, V, Order, Scope) \
+ __atomic_store_n(P, V, Order)
+#define __scoped_atomic_compare_exchange_n(P, E, V, Weak, Success, Failure, \
+ Scope) \
+ __atomic_compare_exchange_n(P, E, V, Weak, Success, Failure)
+#define __scoped_atomic_exchange_n(P, V, Order, Scope) \
+ __atomic_exchange_n(P, V, Order)
+#endif
+
+namespace __csan {
+
+using __sanitizer::u32;
+using __sanitizer::u64;
+using __sanitizer::uptr;
+
+template <u32 MaxAccessSize, u32 CheckAdjacentSlots> class WatchpointTable {
+ static_assert(MaxAccessSize && !(MaxAccessSize & (MaxAccessSize - 1)),
+ "maximum access size must be a power of two");
+
+public:
+ static constexpr u32 SizeBits = __builtin_popcount(MaxAccessSize - 1) + 1;
+ static constexpr u32 AddressBits = 64 - 2 - SizeBits;
+ static constexpr u64 AddressMask = (1ull << AddressBits) - 1;
+ static constexpr u32 OverflowEntries = 2 * CheckAdjacentSlots;
+
+private:
+ // Armed watchpoint (u64):
+ // [63] is_write
+ // [62] consumed = 0
+ // [61:AddressBits] access size
+ // [AddressBits-1:0] address key
+ //
+ // Consumed watchpoint (u64):
+ // [63] peer is_write
+ // [62] consumed = 1
+ // [61:AddressBits] peer access size
+ // [AddressBits-1:0] peer PC
+ static constexpr u64 Invalid = 0;
+ static constexpr u64 ConsumedMask = 1ull << 62;
+ static constexpr u64 WriteMask = 1ull << 63;
+ static constexpr u64 SizeMask = ((1ull << SizeBits) - 1) << AddressBits;
+ static constexpr u32 NumSlots = 1 + 2 * CheckAdjacentSlots;
+
+ u64 *Table;
+
+ static constexpr u64 Encode(u64 Address, u32 Size, bool IsWrite) {
+ return (IsWrite ? WriteMask : 0) | (static_cast<u64>(Size) << AddressBits) |
+ (Address & AddressMask);
+ }
+
+ static constexpr bool Decode(u64 Value, u64 &Address, u32 &Size,
+ bool &IsWrite) {
+ if (Value == Invalid || (Value & ConsumedMask))
+ return false;
+ IsWrite = Value & WriteMask;
+ Size = (Value & SizeMask) >> AddressBits;
+ Address = Value & AddressMask;
+ return true;
+ }
+
+ static constexpr bool Overlaps(u64 A, u32 ASize, u64 B, u32 BSize) {
+ return A < B + BSize && B < A + ASize;
+ }
+
+public:
+ constexpr WatchpointTable(u64 *Table) : Table(Table) {}
+
+ static constexpr u32 EncodeSize(u32 Size) {
+ return Size < MaxAccessSize ? Size : MaxAccessSize;
+ }
+
+ u64 *Find(u64 Key, u32 Size, bool ExpectWrite, u32 Slot, u64 &Encoded) const {
+ for (u32 I = 0; I < NumSlots; ++I) {
+ u64 *Watchpoint = &Table[Slot + I];
+ Encoded = __scoped_atomic_load_n(Watchpoint, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+
+ u64 Address;
+ u32 WatchSize;
+ bool IsWrite;
+ if (!Decode(Encoded, Address, WatchSize, IsWrite))
+ continue;
+ if (ExpectWrite && !IsWrite)
+ continue;
+ if (Overlaps(Address, WatchSize, Key, Size))
+ return Watchpoint;
+ }
+ return nullptr;
+ }
+
+ u64 *Insert(u64 Key, u32 Size, bool IsWrite, u32 Slot) const {
+ const u64 Encoded = Encode(Key, Size, IsWrite);
+ for (u32 I = 0; I < NumSlots; ++I) {
+ u32 Index = Slot + ((I + CheckAdjacentSlots) % NumSlots);
+ u64 *Watchpoint = &Table[Index];
+ u64 Expected = Invalid;
+ if (__scoped_atomic_compare_exchange_n(
+ Watchpoint, &Expected, Encoded, false, __ATOMIC_RELAXED,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE))
+ return Watchpoint;
+ }
+ return nullptr;
+ }
+
+ bool TryConsume(u64 *Watchpoint, u64 Encoded, uptr PC, bool IsWrite,
+ u32 Size) const {
+ u64 Consumed = ConsumedMask | (IsWrite ? WriteMask : 0) |
+ (static_cast<u64>(EncodeSize(Size)) << AddressBits) |
+ (PC & AddressMask);
+ return __scoped_atomic_compare_exchange_n(
+ Watchpoint, &Encoded, Consumed, false, __ATOMIC_RELAXED,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
+ }
+
+ bool Consume(u64 *Watchpoint, void *&Peer, int &PeerAccess,
+ u32 &PeerSize) const {
+ u64 Old = __scoped_atomic_exchange_n(
+ Watchpoint, ConsumedMask, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
+ Peer = reinterpret_cast<void *>(Old & AddressMask);
+ PeerAccess = (Old & WriteMask) ? CSAN_ACCESS_WRITE : 0;
+ PeerSize = (Old & SizeMask) >> AddressBits;
+ return !(Old & ConsumedMask);
+ }
+
+ void Remove(u64 *Watchpoint) const {
+ __scoped_atomic_store_n(Watchpoint, Invalid, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+ }
+};
+
+} // namespace __csan
+
+#ifdef CSAN_DEFINED_SCOPED_ATOMICS
+#undef __scoped_atomic_load_n
+#undef __scoped_atomic_store_n
+#undef __scoped_atomic_compare_exchange_n
+#undef __scoped_atomic_exchange_n
+#undef CSAN_DEFINED_SCOPED_ATOMICS
+#endif
+
+#endif // CSAN_WATCH_H
diff --git a/compiler-rt/lib/csan/offload/CMakeLists.txt b/compiler-rt/lib/csan/offload/CMakeLists.txt
new file mode 100644
index 00000000000000..346be95f20fb96
--- /dev/null
+++ b/compiler-rt/lib/csan/offload/CMakeLists.txt
@@ -0,0 +1,25 @@
+include(FindLibcCommonUtils)
+if(NOT OS_NAME MATCHES "Linux" OR NOT TARGET llvm-libc-common-utilities)
+ return()
+endif()
+
+set(CSAN_OFFLOAD_SOURCES
+ csan_offload_hsa_interceptors.cpp
+ csan_offload_report.cpp
+ )
+
+set(CSAN_OFFLOAD_HEADERS
+ csan_offload.h
+ ../csan_defs.h
+ ../csan_offload_packet.h
+ )
+
+add_compiler_rt_runtime(clang_rt.csan_offload
+ STATIC
+ ARCHS ${UBSAN_SUPPORTED_ARCH}
+ SOURCES ${CSAN_OFFLOAD_SOURCES}
+ ADDITIONAL_HEADERS ${CSAN_OFFLOAD_HEADERS}
+ OBJECT_LIBS RTSanitizerOffload
+ CFLAGS ${CSAN_CFLAGS}
+ LINK_LIBS llvm-libc-common-utilities
+ PARENT_TARGET csan)
diff --git a/compiler-rt/lib/csan/offload/csan_offload.h b/compiler-rt/lib/csan/offload/csan_offload.h
new file mode 100644
index 00000000000000..f37d39a37162e5
--- /dev/null
+++ b/compiler-rt/lib/csan/offload/csan_offload.h
@@ -0,0 +1,31 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Internal declarations for host-side ConcurrencySanitizer offload reporting.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef CSAN_OFFLOAD_H
+#define CSAN_OFFLOAD_H
+
+#include "csan_defs.h"
+#include "csan_offload_packet.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+
+namespace __csan {
+
+u32 HandleOffloadReport(void *Port, u32 Lanes);
+
+} // namespace __csan
+
+extern "C" {
+SANITIZER_INTERFACE_ATTRIBUTE void __csan_offload_init();
+}
+
+#endif // CSAN_OFFLOAD_H
diff --git a/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp b/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
new file mode 100644
index 00000000000000..c08dc58a49637d
--- /dev/null
+++ b/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
@@ -0,0 +1,370 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// HSA interceptors for host-side ConcurrencySanitizer offload support.
+///
+//===----------------------------------------------------------------------===//
+
+#include <dlfcn.h>
+
+#include "csan_offload.h"
+#include "interception/interception.h"
+#include "sanitizer_common/sanitizer_atomic.h"
+#include "sanitizer_common/sanitizer_common.h"
+#include "sanitizer_common/sanitizer_flag_parser.h"
+#include "sanitizer_common/sanitizer_flags.h"
+#include "sanitizer_common/sanitizer_libc.h"
+#include "sanitizer_common/sanitizer_mutex.h"
+#include "sanitizer_common/sanitizer_offload.h"
+#include "sanitizer_common/sanitizer_platform.h"
+#include "sanitizer_common/sanitizer_symbolizer.h"
+#include "sanitizer_common/sanitizer_vector.h"
+
+#if !SANITIZER_LINUX
+#error "Offload CSan reporting is supported on Linux only"
+#endif
+
+#if SANITIZER_GLIBC
+#pragma weak dlvsym
+#endif
+
+using namespace __sanitizer;
+using namespace __csan;
+
+namespace {
+
+static StaticSpinMutex InitMutex;
+static StaticSpinMutex HsaMutex;
+static atomic_uint8_t Initialized;
+static void *HsaHandle;
+
+struct DeviceAllocation {
+ hsa_agent_t Agent;
+ void *Ptr;
+};
+
+Mutex AllocationMutex;
+InternalMmapVectorNoCtor<DeviceAllocation> Allocations;
+uptr HsaRefs;
+
+void Initialize() {
+ if (LIKELY(atomic_load(&Initialized, memory_order_acquire)))
+ return;
+ SpinMutexLock L(&InitMutex);
+ if (atomic_load(&Initialized, memory_order_relaxed))
+ return;
+ SanitizerToolName = "ConcurrencySanitizer";
+ CacheBinaryName();
+ SetCommonFlagsDefaults();
+ {
+ CommonFlags cf;
+ cf.CopyFrom(*common_flags());
+ if (const char *Path = GetEnv("CSAN_SYMBOLIZER_PATH"))
+ cf.external_symbolizer_path = Path;
+ cf.stack_trace_format = " #%n %f %S (%p)";
+ OverrideCommonFlags(cf);
+ }
+ FlagParser Parser;
+ RegisterCommonFlags(&Parser);
+ Parser.ParseStringFromEnv("CSAN_OPTIONS");
+ InitializeCommonFlags();
+ Offload::Get().RegisterHandler(HandleOffloadReport);
+ Atexit([] { Offload::Get().UntrackImages(); });
+ AddDieCallback([] { Offload::Get().UntrackImages(); });
+
+ // Mark ready before LateInitialize, as it can be reentrant through dlsym.
+ atomic_store(&Initialized, 1, memory_order_release);
+ Symbolizer::LateInitialize();
+}
+
+} // namespace
+
+static void BindRealDlsym();
+static void *HsaSymbol(const char *Name);
+
+#define CSAN_HSA_ENTER(name) \
+ Initialize(); \
+ if (UNLIKELY(!REAL(name) || REAL(name) == name)) { \
+ REAL(name) = reinterpret_cast<decltype(REAL(name))>(HsaSymbol(#name)); \
+ if (UNLIKELY(!REAL(name) || REAL(name) == name)) { \
+ Report("ERROR: %s: cannot find %s in this process\n", SanitizerToolName, \
+ #name); \
+ Die(); \
+ } \
+ }
+
+#define CSAN_HSA_FORWARD(name, ...) \
+ CSAN_HSA_ENTER(name); \
+ 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
+#define CSAN_HSA_WRAPS(X) \
+ X(hsa_init) \
+ X(hsa_shut_down) \
+ X(hsa_executable_freeze) \
+ X(hsa_executable_destroy)
+
+static void *WrapperFor(const char *Name) {
+#define CSAN_HSA_WRAP(Fn) \
+ if (!internal_strcmp(Name, #Fn)) \
+ return reinterpret_cast<void *>(Fn);
+ CSAN_HSA_WRAPS(CSAN_HSA_WRAP)
+#undef CSAN_HSA_WRAP
+ return nullptr;
+}
+
+static bool FromHsa(void *P) {
+ Dl_info Info = {};
+ if (!dladdr(P, &Info) || !Info.dli_fname)
+ return false;
+ return internal_strstr(Info.dli_fname, SANITIZER_HSA_LIBRARY);
+}
+
+// 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]]
+ return REAL(dlsym)(Handle, Name);
+
+ void *Sym = REAL(dlsym)(Handle, Name);
+ if (!Sym || !Name)
+ return Sym;
+
+ void *Wrapper = WrapperFor(Name);
+ if (!Wrapper || !FromHsa(Sym))
+ return Sym;
+ return Wrapper;
+}
+#else
+DEFINE_REAL(void *, dlsym, void *, const char *)
+#endif
+
+static void BindRealDlsym() {
+ if (LIKELY(REAL(dlsym)))
+ return;
+#if SANITIZER_GLIBC
+ static const char *kVers[] = {"GLIBC_2.34", "GLIBC_2.17", "GLIBC_2.2.5",
+ "GLIBC_2.0"};
+ if (dlvsym) {
+ for (const char *Ver : kVers) {
+ if (void *P = dlvsym(RTLD_NEXT, "dlsym", Ver)) {
+ REAL(dlsym) = reinterpret_cast<decltype(REAL(dlsym))>(P);
+ return;
+ }
+ }
+ }
+#endif
+ Report("ERROR: %s: cannot bind dlsym\n", SanitizerToolName);
+ Die();
+}
+
+static void *HsaSymbol(const char *Name) {
+ BindRealDlsym();
+ if (!HsaHandle) {
+ SpinMutexLock L(&HsaMutex);
+ if (HsaHandle)
+ return REAL(dlsym)(HsaHandle, Name);
+ constexpr const char *Names[] = {"libhsa-runtime64.so.1",
+ "libhsa-runtime64.so"};
+ for (const char *Name : Names)
+ if (void *H = dlopen(Name, RTLD_LAZY | RTLD_NOLOAD))
+ HsaHandle = H;
+ for (const char *Name : Names)
+ if (!HsaHandle)
+ HsaHandle = dlopen(Name, RTLD_LAZY | RTLD_LOCAL);
+ }
+ return HsaHandle ? REAL(dlsym)(HsaHandle, Name) : nullptr;
+}
+
+template <typename T> static T HsaFunction(const char *Name) {
+ return reinterpret_cast<T>(HsaSymbol(Name));
+}
+
+static bool Lookup(hsa_executable_t Executable, const char *Name,
+ hsa_agent_t Agent, u64 *Addr) {
+ auto GetSymbol = HsaFunction<decltype(&hsa_executable_get_symbol_by_name)>(
+ "hsa_executable_get_symbol_by_name");
+ auto GetInfo = HsaFunction<decltype(&hsa_executable_symbol_get_info)>(
+ "hsa_executable_symbol_get_info");
+ if (!GetSymbol || !GetInfo)
+ return false;
+
+ hsa_executable_symbol_t Symbol;
+ if (GetSymbol(Executable, Name, &Agent, &Symbol) != HSA_STATUS_SUCCESS)
+ return false;
+ *Addr = 0;
+ return GetInfo(Symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, Addr) ==
+ HSA_STATUS_SUCCESS &&
+ *Addr;
+}
+
+static DeviceAllocation *FindAllocation(hsa_agent_t Agent) {
+ for (uptr I = 0; I < Allocations.size(); ++I)
+ if (Allocations[I].Agent.handle == Agent.handle)
+ return &Allocations[I];
+ return nullptr;
+}
+
+struct BindContext {
+ hsa_executable_t Executable;
+ bool Found;
+ bool Success;
+};
+
+// The watchpoint table is shared between all executables loaded on the device.
+static hsa_status_t BindAgent(hsa_agent_t Agent, void *Data) {
+ BindContext &Ctx = *reinterpret_cast<BindContext *>(Data);
+ auto AgentInfo =
+ HsaFunction<decltype(&hsa_agent_get_info)>("hsa_agent_get_info");
+ auto Copy = HsaFunction<decltype(&hsa_memory_copy)>("hsa_memory_copy");
+ auto Free = HsaFunction<decltype(&hsa_amd_memory_pool_free)>(
+ "hsa_amd_memory_pool_free");
+ if (!AgentInfo || !Copy || !Free) {
+ Ctx.Success = false;
+ return HSA_STATUS_SUCCESS;
+ }
+
+ hsa_device_type_t Type;
+ if (AgentInfo(Agent, HSA_AGENT_INFO_DEVICE, &Type) != HSA_STATUS_SUCCESS ||
+ Type != HSA_DEVICE_TYPE_GPU)
+ return HSA_STATUS_SUCCESS;
+
+ u64 PointerAddr = 0;
+ if (!Lookup(Ctx.Executable, "__csan_watchpoint_table", Agent, &PointerAddr))
+ return HSA_STATUS_SUCCESS;
+ Ctx.Found = true;
+
+ // Create an allocation for the table if one does not already exist.
+ DeviceAllocation *Allocation = FindAllocation(Agent);
+ if (!Allocation) {
+ hsa_amd_memory_pool_t Pool;
+ void *Ptr = nullptr;
+ if (!Offload::Get().GetMemoryPool(Agent, &Pool) ||
+ !Offload::Get().Allocate(Pool, CSAN_WATCHPOINT_TABLE_BYTES, &Ptr)) {
+ Ctx.Success = false;
+ return HSA_STATUS_SUCCESS;
+ }
+
+ // Fresh pages from MMap are always zero initialized, as required.
+ void *Zeros =
+ MmapOrDie(CSAN_WATCHPOINT_TABLE_BYTES, "CSan watchpoint table");
+ bool Copied =
+ Copy(Ptr, Zeros, CSAN_WATCHPOINT_TABLE_BYTES) == HSA_STATUS_SUCCESS;
+ UnmapOrDie(Zeros, CSAN_WATCHPOINT_TABLE_BYTES);
+ if (!Copied) {
+ Free(Ptr);
+ Ctx.Success = false;
+ return HSA_STATUS_SUCCESS;
+ }
+ Allocations.push_back({Agent, Ptr});
+ Allocation = &Allocations.back();
+ }
+
+ // Update the pointer on the device to the associated allocation.
+ if (Copy(reinterpret_cast<void *>(PointerAddr), &Allocation->Ptr,
+ sizeof(Allocation->Ptr)) != HSA_STATUS_SUCCESS)
+ Ctx.Success = false;
+ return HSA_STATUS_SUCCESS;
+}
+
+static void BindWatchpointTable(hsa_executable_t Executable) {
+ auto Iterate =
+ HsaFunction<decltype(&hsa_iterate_agents)>("hsa_iterate_agents");
+ if (!Iterate)
+ return;
+
+ // Check if we need to set up the watchpoint table.
+ Lock L(&AllocationMutex);
+ BindContext Ctx = {Executable, false, true};
+ if (Iterate(BindAgent, &Ctx) != HSA_STATUS_SUCCESS)
+ Ctx.Success = false;
+ if (Ctx.Found && !Ctx.Success) {
+ Report("ERROR: %s: cannot initialize device watchpoint table\n",
+ SanitizerToolName);
+ Die();
+ }
+}
+
+static void RetainAllocations() {
+ Lock L(&AllocationMutex);
+ ++HsaRefs;
+}
+
+static void ReleaseAllocations() {
+ Lock L(&AllocationMutex);
+ if (!HsaRefs || --HsaRefs)
+ return;
+ auto Free = HsaFunction<decltype(&hsa_amd_memory_pool_free)>(
+ "hsa_amd_memory_pool_free");
+ if (Free)
+ for (uptr I = 0; I < Allocations.size(); ++I)
+ Free(Allocations[I].Ptr);
+ Allocations.clear();
+}
+
+INTERCEPTOR(hsa_status_t, hsa_init, void) {
+ CSAN_HSA_ENTER(hsa_init);
+
+ hsa_status_t Status = REAL(hsa_init)();
+ if (Status != HSA_STATUS_SUCCESS)
+ return Status;
+
+ if (!Offload::Get().Init()) {
+ Report("ERROR: %s: cannot initialize HSA offload support\n",
+ SanitizerToolName);
+ Die();
+ }
+ RetainAllocations();
+ return Status;
+}
+
+INTERCEPTOR(hsa_status_t, hsa_shut_down, void) {
+ CSAN_HSA_ENTER(hsa_shut_down);
+
+ ReleaseAllocations();
+ Offload::Get().Shutdown();
+ return REAL(hsa_shut_down)();
+}
+
+INTERCEPTOR(hsa_status_t, hsa_executable_freeze, hsa_executable_t Executable,
+ const char *Options) {
+ CSAN_HSA_FORWARD(hsa_executable_freeze, Executable, Options);
+
+ hsa_status_t Status = REAL(hsa_executable_freeze)(Executable, Options);
+ if (Status == HSA_STATUS_SUCCESS) {
+ BindWatchpointTable(Executable);
+ Offload::Get().TrackExecutable(Executable);
+ }
+ return Status;
+}
+
+INTERCEPTOR(hsa_status_t, hsa_executable_destroy, hsa_executable_t Executable) {
+ CSAN_HSA_FORWARD(hsa_executable_destroy, Executable);
+
+ Offload::Get().UntrackExecutable(Executable);
+ return REAL(hsa_executable_destroy)(Executable);
+}
+
+extern "C" void __csan_offload_init() { Initialize(); }
+
+#if SANITIZER_CAN_USE_PREINIT_ARRAY
+__attribute__((section(".preinit_array"), used)) static void (
+ *csan_offload_preinit)(void) = __csan_offload_init;
+#endif
+
+__attribute__((constructor(0))) static void CsanOffloadDynInit() {
+ __csan_offload_init();
+}
diff --git a/compiler-rt/lib/csan/offload/csan_offload_report.cpp b/compiler-rt/lib/csan/offload/csan_offload_report.cpp
new file mode 100644
index 00000000000000..1c477d33d1037d
--- /dev/null
+++ b/compiler-rt/lib/csan/offload/csan_offload_report.cpp
@@ -0,0 +1,184 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Host-side device race report generation.
+///
+//===----------------------------------------------------------------------===//
+
+#include "csan_offload.h"
+
+#include "sanitizer_common/sanitizer_common.h"
+#include "sanitizer_common/sanitizer_flags.h"
+#include "sanitizer_common/sanitizer_libc.h"
+#include "sanitizer_common/sanitizer_mutex.h"
+#include "sanitizer_common/sanitizer_offload.h"
+#include "sanitizer_common/sanitizer_report_decorator.h"
+#include "sanitizer_common/sanitizer_stacktrace_printer.h"
+#include "sanitizer_common/sanitizer_symbolizer.h"
+#include "shared/rpc.h"
+
+using namespace __sanitizer;
+
+namespace __csan {
+namespace {
+
+class Decorator : public SanitizerCommonDecorator {
+public:
+ const char *Access() { return Blue(); }
+ const char *Location() { return Green(); }
+};
+
+const char *KindName(u32 Kind) {
+ switch (Kind) {
+ case CSAN_RACE_UNKNOWN_ORIGIN:
+ return "data race of unknown origin";
+ case CSAN_RACE_INTRA_WAVE:
+ return "intra-wave data race";
+ default:
+ return "data race";
+ }
+}
+
+const char *MopDesc(bool First, u32 Type) {
+ const bool Write = Type & CSAN_ACCESS_WRITE;
+ return First ? (Write ? "Write" : "Read")
+ : (Write ? "Previous write" : "Previous read");
+}
+
+// A pair of PCs that have already been reported as racing. The device samples
+// the same static access from thousands of threads, so without this every
+// launch buries the user in duplicates of one bug.
+struct RacyPcs {
+ u64 pc[2];
+
+ bool operator==(const RacyPcs &other) const {
+ if (pc[0] == other.pc[0] && pc[1] == other.pc[1])
+ return true;
+ return pc[0] == other.pc[1] && pc[1] == other.pc[0];
+ }
+};
+
+Mutex RacyMutex;
+InternalMmapVectorNoCtor<RacyPcs> RacyPcsSeen;
+
+bool FindRacyPcs(const RacyPcs &Racy) {
+ for (uptr I = 0; I < RacyPcsSeen.size(); ++I) {
+ if (Racy == RacyPcsSeen[I]) {
+ VReport(2, "%s: suppressing report as doubled\n", SanitizerToolName);
+ return true;
+ }
+ }
+ return false;
+}
+
+bool HandleRacyPcs(const __csan_gpu_race &R) {
+ RacyPcs Racy = {{R.pc, R.peer_pc}};
+ {
+ ReadLock L(&RacyMutex);
+ if (FindRacyPcs(Racy))
+ return true;
+ }
+ Lock L(&RacyMutex);
+ if (FindRacyPcs(Racy))
+ return true;
+ RacyPcsSeen.push_back(Racy);
+ return false;
+}
+
+void PrintFrames(SymbolizedStack *Frames, u64 PC) {
+ if (!Frames) {
+ Printf(" #0 (%p)\n", (void *)(uptr)PC);
+ return;
+ }
+ const SymbolizedStack *F = SkipInternalFrames(Frames);
+ if (!F)
+ F = Frames;
+ int N = 0;
+ for (; F; F = F->next, ++N) {
+ InternalScopedString Res;
+ StackTracePrinter::GetOrInit()->RenderFrame(
+ &Res, common_flags()->stack_trace_format, N, F->info.address, &F->info,
+ common_flags()->symbolize_vs_style, common_flags()->strip_path_prefix);
+ Printf("%s\n", Res.data());
+ }
+}
+
+} // namespace
+
+void PrintOffloadReport(const __csan_gpu_race &R) {
+ if (HandleRacyPcs(R))
+ return;
+
+ Decorator D;
+ Printf("==================\n");
+ Printf("%s", D.Warning());
+ Printf("WARNING: ConcurrencySanitizer: %s\n", KindName(R.kind));
+ Printf("%s", D.Default());
+
+ Printf("%s", D.Access());
+ Printf(" %s of size %u at 0x%zx in block (%u,%u,%u) thread (%u,%u,%u) "
+ "lane %u:\n",
+ MopDesc(true, R.access_type), R.size, (uptr)R.addr, R.block[0],
+ R.block[1], R.block[2], R.thread[0], R.thread[1], R.thread[2], R.lane);
+ Printf("%s", D.Default());
+ SymbolizedStack *This = Offload::Get().Symbolize((uptr)R.pc);
+ PrintFrames(This, R.pc);
+
+ Printf("%s", D.Access());
+ if (R.peer_pc) {
+ Printf(" %s of size %u at 0x%zx:\n", MopDesc(false, R.peer_access_type),
+ R.peer_size, (uptr)R.addr);
+ Printf("%s", D.Default());
+ SymbolizedStack *Peer = Offload::Get().Symbolize((uptr)R.peer_pc);
+ PrintFrames(Peer, R.peer_pc);
+ if (Peer)
+ Peer->ClearAll();
+ } else if (R.kind == CSAN_RACE_INTRA_WAVE) {
+ Printf(" Previous access by lane %u in the same wave\n", R.peer_lane);
+ Printf("%s", D.Default());
+ } else {
+ Printf(" Previous access of unknown origin\n");
+ Printf("%s", D.Default());
+ }
+
+ DataInfo Loc;
+ if (Offload::Get().SymbolizeData((uptr)R.addr, &Loc)) {
+ Printf("%s", D.Location());
+ if (Loc.size)
+ Printf(" Location is global '%s' of size %zu at 0x%zx\n", Loc.name,
+ Loc.size, (uptr)R.addr);
+ else
+ Printf(" Location is global '%s' at 0x%zx\n", Loc.name, (uptr)R.addr);
+ Printf("%s", D.Default());
+ Loc.Clear();
+ }
+
+ if (This) {
+ const SymbolizedStack *User = SkipInternalFrames(This);
+ ReportErrorSummary(KindName(R.kind), (User ? User : This)->info);
+ This->ClearAll();
+ }
+
+ Printf("==================\n");
+}
+
+u32 HandleOffloadReport(void *PortPtr, u32) {
+ auto &Port = *reinterpret_cast<rpc::Server::Port *>(PortPtr);
+ if (Port.get_opcode() != SANITIZER_OFFLOAD_CSAN)
+ return rpc::RPC_UNHANDLED_OPCODE;
+
+ Port.recv([&](rpc::Buffer *Buffer, u32) {
+ __csan_gpu_race R;
+ internal_memcpy(&R, Buffer->data, sizeof(R));
+ PrintOffloadReport(R);
+ });
+ return rpc::RPC_SUCCESS;
+}
+
+} // namespace __csan
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h b/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
index b8a6ac95d607b2..f9e02938fc437b 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
@@ -489,6 +489,9 @@ using namespace __sanitizer;
namespace __ubsan {
using namespace __sanitizer;
}
+namespace __csan {
+using namespace __sanitizer;
+}
namespace __xray {
using namespace __sanitizer;
}
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_offload_opcodes.h b/compiler-rt/lib/sanitizer_common/sanitizer_offload_opcodes.h
index 4f3d0ce8e9fd23..6285e4d527face 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_offload_opcodes.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_offload_opcodes.h
@@ -20,6 +20,7 @@
enum {
SANITIZER_OFFLOAD_UBSAN = SANITIZER_OFFLOAD_OPCODE(0),
+ SANITIZER_OFFLOAD_CSAN = SANITIZER_OFFLOAD_OPCODE(1),
};
#undef SANITIZER_OFFLOAD_RPC_BASE
diff --git a/compiler-rt/test/csan/AMDGPU/aba-race.hip b/compiler-rt/test/csan/AMDGPU/aba-race.hip
new file mode 100644
index 00000000000000..e5da74ff7b7639
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/aba-race.hip
@@ -0,0 +1,33 @@
+// RUN: %clang_hip_csan -O2 %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile int Global;
+
+__global__ void kernel() {
+ unsigned Id = CSAN_TID();
+ int Sink = 0;
+ RACE_UNTIL_FOUND(i) {
+ if (Id & 1) {
+ Global = 1;
+ Global = 0;
+ } else {
+ Sink += Global;
+ }
+ }
+ if (Sink == -1)
+ __builtin_trap();
+}
+
+int main() {
+ kernel<<<64, 1>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: {{Read|Write}} of size 4
+// CHECK: Previous {{read|write}} of size 4
+// CHECK: #0 {{.*}}aba-race.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}Global{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/AMDGPU/array-race.hip b/compiler-rt/test/csan/AMDGPU/array-race.hip
new file mode 100644
index 00000000000000..052befa554d99b
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/array-race.hip
@@ -0,0 +1,22 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile int data[64];
+
+__global__ void kernel() {
+ unsigned Slot = CSAN_TID() % 64;
+ RACE_UNTIL_FOUND(i)
+ data[Slot]++;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}array-race.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}data{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/AMDGPU/atomic-nonatomic.hip b/compiler-rt/test/csan/AMDGPU/atomic-nonatomic.hip
new file mode 100644
index 00000000000000..729188ef07b8d6
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/atomic-nonatomic.hip
@@ -0,0 +1,26 @@
+// RUN: %clang_hip_csan -O2 %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ int Global;
+
+__global__ void kernel() {
+ unsigned Id = CSAN_TID();
+ RACE_UNTIL_FOUND(i) {
+ if (Id & 1)
+ __atomic_fetch_add(&Global, 1, __ATOMIC_RELAXED);
+ else
+ Global = i;
+ }
+}
+
+int main() {
+ kernel<<<64, 1>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}atomic-nonatomic.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}Global{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/AMDGPU/atomic.hip b/compiler-rt/test/csan/AMDGPU/atomic.hip
new file mode 100644
index 00000000000000..5a9395dc8059c9
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/atomic.hip
@@ -0,0 +1,17 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --allow-empty
+
+__device__ int global;
+
+__global__ void kernel() {
+ for (int I = 0; I < 1024; ++I)
+ __atomic_fetch_add(&global, 1, __ATOMIC_RELAXED);
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/AMDGPU/disjoint.hip b/compiler-rt/test/csan/AMDGPU/disjoint.hip
new file mode 100644
index 00000000000000..b8399f59a8463e
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/disjoint.hip
@@ -0,0 +1,20 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --allow-empty
+
+#include "race.h"
+
+__device__ int data[64 * 64];
+
+__global__ void kernel() {
+ unsigned Id = CSAN_TID();
+ for (int I = 0; I < 1024; ++I)
+ data[Id % (64 * 64)] += I;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/AMDGPU/global-race.hip b/compiler-rt/test/csan/AMDGPU/global-race.hip
new file mode 100644
index 00000000000000..90970929800f0f
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/global-race.hip
@@ -0,0 +1,23 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+// RUN: %clang_hip_csan -O2 %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile int global;
+
+__global__ void kernel() {
+ RACE_UNTIL_FOUND(i)
+ global++;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}global-race.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}global{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/AMDGPU/helper-race.hip b/compiler-rt/test/csan/AMDGPU/helper-race.hip
new file mode 100644
index 00000000000000..f61c750a8c1e31
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/helper-race.hip
@@ -0,0 +1,26 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+// RUN: %clang_hip_csan -O2 %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile int global;
+
+__device__ __attribute__((always_inline)) void inc() { global++; }
+
+__global__ void kernel() {
+ RACE_UNTIL_FOUND(i)
+ inc();
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}inc{{.*}}helper-race.hip:{{[0-9]+}}
+// CHECK: #1 {{.*}}kernel{{.*}}helper-race.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}global{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/AMDGPU/lds-disjoint.hip b/compiler-rt/test/csan/AMDGPU/lds-disjoint.hip
new file mode 100644
index 00000000000000..d3dcae7a34427b
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/lds-disjoint.hip
@@ -0,0 +1,21 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --allow-empty
+
+#include "race.h"
+
+// Each workgroup has its own LDS. Many groups touching the same offset must
+// not look like a race once the key carries the workgroup index.
+__global__ void kernel() {
+ __shared__ int Shared[64];
+ unsigned Slot = CSAN_TID() % 64;
+ for (int I = 0; I < 1024; ++I)
+ Shared[Slot] += I;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/AMDGPU/lds-race.hip b/compiler-rt/test/csan/AMDGPU/lds-race.hip
new file mode 100644
index 00000000000000..33cf99a2690cd8
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/lds-race.hip
@@ -0,0 +1,21 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__global__ void kernel() {
+ __shared__ volatile int Shared[64];
+ RACE_UNTIL_FOUND(i)
+ Shared[0] += i;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: {{Write|Read}} of size 4 at 0x{{[1-9a-f][0-9a-f]+}}
+// CHECK: #0 {{.*}}lds-race.hip:{{[0-9]+}}
+// CHECK-NOT: Location is global
diff --git a/compiler-rt/test/csan/AMDGPU/lit.local.cfg.py b/compiler-rt/test/csan/AMDGPU/lit.local.cfg.py
new file mode 100644
index 00000000000000..e4dd4ea03d4d05
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/lit.local.cfg.py
@@ -0,0 +1,11 @@
+# Discover only tests for GPU frontends that are usable on this system.
+config.suffixes = []
+if "csan-hip" in config.available_features:
+ config.suffixes.append(".hip")
+if "csan-openmp-offload" in config.available_features:
+ config.suffixes.append(".cpp")
+
+if not config.suffixes:
+ config.unsupported = True
+else:
+ config.parallelism_group = "gpu"
diff --git a/compiler-rt/test/csan/AMDGPU/memcpy-large-race.hip b/compiler-rt/test/csan/AMDGPU/memcpy-large-race.hip
new file mode 100644
index 00000000000000..93f2bce077f90c
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/memcpy-large-race.hip
@@ -0,0 +1,28 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+#define N 512
+
+__device__ int Dst[N];
+__device__ int Src[N];
+
+__global__ void kernel() {
+ unsigned Id = CSAN_TID();
+ RACE_UNTIL_FOUND(i) {
+ if (Id == 0)
+ __builtin_memcpy((void *)Dst, (const void *)Src, sizeof(Dst));
+ else
+ Dst[N - 1] = Id;
+ }
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}memcpy-large-race.hip:{{[0-9]+}}
diff --git a/compiler-rt/test/csan/AMDGPU/memcpy-race.hip b/compiler-rt/test/csan/AMDGPU/memcpy-race.hip
new file mode 100644
index 00000000000000..b0b504cf2ead9b
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/memcpy-race.hip
@@ -0,0 +1,22 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile unsigned Len = 32;
+__device__ int Dst[64];
+__device__ int Src[64];
+
+__global__ void kernel() {
+ RACE_UNTIL_FOUND(i)
+ __builtin_memcpy((void *)Dst, (const void *)Src, Len);
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}memcpy-race.hip:{{[0-9]+}}
diff --git a/compiler-rt/test/csan/AMDGPU/memmove-race.hip b/compiler-rt/test/csan/AMDGPU/memmove-race.hip
new file mode 100644
index 00000000000000..2807a1c625dbec
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/memmove-race.hip
@@ -0,0 +1,28 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+#define N 512
+
+__device__ int Buf[N];
+
+__global__ void kernel() {
+ unsigned Id = CSAN_TID();
+ RACE_UNTIL_FOUND(i) {
+ if (Id == 0)
+ __builtin_memmove((void *)Buf, (const void *)(Buf + 1),
+ (N - 1) * sizeof(int));
+ else
+ Buf[N - 1] = Id;
+ }
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}memmove-race.hip:{{[0-9]+}}
diff --git a/compiler-rt/test/csan/AMDGPU/openmp-race.cpp b/compiler-rt/test/csan/AMDGPU/openmp-race.cpp
new file mode 100644
index 00000000000000..2c053c0e8e6681
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/openmp-race.cpp
@@ -0,0 +1,19 @@
+// RUN: %clang_omp_offload_csan -fno-exceptions %s -o %t
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+int main() {
+ int X = 0;
+#pragma omp target teams num_teams(1) thread_limit(64) map(tofrom : X)
+#pragma omp parallel num_threads(64)
+ {
+ volatile int *P = (volatile int *)&X;
+ RACE_UNTIL_FOUND(I)
+ *P = I;
+ }
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}openmp-race.cpp:{{[0-9]+}}
diff --git a/compiler-rt/test/csan/AMDGPU/race.h b/compiler-rt/test/csan/AMDGPU/race.h
new file mode 100644
index 00000000000000..997c23f09aa3f1
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/race.h
@@ -0,0 +1,32 @@
+#ifndef CSAN_TEST_RACE_H
+#define CSAN_TEST_RACE_H
+
+#if defined(__HIP__)
+extern "C" __device__ unsigned long long __csan_get_num_data_races(void);
+# define CSAN_DEVICE __device__
+#elif defined(__AMDGCN__)
+extern "C" unsigned long long __csan_get_num_data_races(void);
+# define CSAN_DEVICE
+#else
+extern "C" unsigned long long __csan_get_num_data_races(void);
+# define CSAN_DEVICE
+#endif
+
+CSAN_DEVICE static inline int race_found(void) noexcept {
+ return __csan_get_num_data_races() != 0;
+}
+
+#define RACE_MAX_ITERS (1 << 20)
+#define RACE_UNTIL_FOUND(i) \
+ for (int i = 0; i < RACE_MAX_ITERS && !race_found(); ++i)
+
+#if defined(__HIP_DEVICE_COMPILE__)
+# include <gpuintrin.h>
+# define CSAN_TID() \
+ (__gpu_num_threads(__GPU_X_DIM) * __gpu_block_id(__GPU_X_DIM) + \
+ __gpu_thread_id(__GPU_X_DIM))
+#else
+# define CSAN_TID() 0u
+#endif
+
+#endif
diff --git a/compiler-rt/test/csan/AMDGPU/shared-watchpoints.hip b/compiler-rt/test/csan/AMDGPU/shared-watchpoints.hip
new file mode 100644
index 00000000000000..e19733a3466158
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/shared-watchpoints.hip
@@ -0,0 +1,59 @@
+// RUN: %clang_hip_csan -O2 -fPIC -shared -Xarch_host \
+// RUN: -fno-sanitize=concurrency -DLIB_A %s -o %t.a.so %hip_libs
+// RUN: %clang_hip_csan -O2 -fPIC -shared -Xarch_host \
+// RUN: -fno-sanitize=concurrency -DLIB_B %s -o %t.b.so %hip_libs
+// RUN: %clangxx -x c++ %s -x none %t.a.so %t.b.so \
+// RUN: -Wl,-rpath -Wl,%T -o %t
+// RUN: %run %t | FileCheck %s
+
+#if defined(LIB_A) || defined(LIB_B)
+
+#include <stdint.h>
+
+extern "C" __device__ uint64_t *__csan_watchpoint_table;
+extern "C" int hipMemcpy(void *, const void *, unsigned long, int);
+static constexpr int hipMemcpyDeviceToHost = 2;
+
+#if defined(LIB_A)
+#define PRINT_TABLE print_table_a
+#define TABLE_LABEL "TABLE_A"
+#else
+#define PRINT_TABLE print_table_b
+#define TABLE_LABEL "TABLE_B"
+#endif
+
+__global__ void read_table(uint64_t *out) {
+ *out = reinterpret_cast<uint64_t>(__csan_watchpoint_table);
+}
+
+extern "C" uint64_t PRINT_TABLE() {
+ uint64_t *device_out;
+ uint64_t result = 0;
+ if (hipMalloc(reinterpret_cast<void **>(&device_out), sizeof(result)))
+ return 0;
+ read_table<<<1, 1>>>(device_out);
+ if (hipMemcpy(&result, device_out, sizeof(result), hipMemcpyDeviceToHost))
+ result = 0;
+ hipFree(device_out);
+ return result;
+}
+
+#else
+
+#include <stdint.h>
+#include <stdio.h>
+
+extern "C" uint64_t print_table_a();
+extern "C" uint64_t print_table_b();
+
+int main() {
+ uint64_t a = print_table_a();
+ uint64_t b = print_table_b();
+ printf("TABLE: %#llx %#llx\n", static_cast<unsigned long long>(a),
+ static_cast<unsigned long long>(b));
+ return !a || a != b;
+}
+
+// CHECK: TABLE: [[TABLE:0x[0-9a-f]+]] [[TABLE]]
+
+#endif
diff --git a/compiler-rt/test/csan/AMDGPU/single-thread.hip b/compiler-rt/test/csan/AMDGPU/single-thread.hip
new file mode 100644
index 00000000000000..a35d77bc76033b
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/single-thread.hip
@@ -0,0 +1,17 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --allow-empty
+
+__device__ int Global;
+
+__global__ void kernel() {
+ for (int I = 0; I < 1024; ++I)
+ Global++;
+}
+
+int main() {
+ kernel<<<1, 1>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/AMDGPU/write-read-race.hip b/compiler-rt/test/csan/AMDGPU/write-read-race.hip
new file mode 100644
index 00000000000000..c63d8538ea5c88
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/write-read-race.hip
@@ -0,0 +1,26 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+__device__ volatile int Data;
+
+__global__ void kernel() {
+ int Sink = 0;
+ RACE_UNTIL_FOUND(i) {
+ Data = i;
+ Sink += Data;
+ }
+ if (Sink == -1)
+ __builtin_trap();
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: {{.*}}race
+// CHECK: #0 {{.*}}write-read-race.hip:{{[0-9]+}}
+// CHECK: Location is global '{{.*}}Data{{.*}}' of size {{[0-9]+}} at 0x{{.*}}
diff --git a/compiler-rt/test/csan/CMakeLists.txt b/compiler-rt/test/csan/CMakeLists.txt
new file mode 100644
index 00000000000000..6df0f1f1de3d11
--- /dev/null
+++ b/compiler-rt/test/csan/CMakeLists.txt
@@ -0,0 +1,25 @@
+set(CSAN_LIT_TESTS_DIR ${CMAKE_CURRENT_SOURCE_DIR})
+
+# Device compiler-rt only builds the GPU archive, the tests run from the host.
+if(COMPILER_RT_GPU_BUILD)
+ return()
+endif()
+
+set(CSAN_TESTSUITES)
+set(CSAN_TEST_DEPS ${SANITIZER_COMMON_LIT_TEST_DEPS})
+list(APPEND CSAN_TEST_DEPS csan)
+
+set(CSAN_TEST_ARCH ${UBSAN_SUPPORTED_ARCH})
+foreach(arch ${CSAN_TEST_ARCH})
+ set(CSAN_TEST_TARGET_ARCH ${arch})
+ get_test_cc_for_arch(${arch} CSAN_TEST_TARGET_CC CSAN_TEST_TARGET_CFLAGS)
+ set(CONFIG_NAME ${arch})
+ configure_lit_site_cfg(
+ ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
+ ${CMAKE_CURRENT_BINARY_DIR}/${CONFIG_NAME}/lit.site.cfg.py)
+ list(APPEND CSAN_TESTSUITES ${CMAKE_CURRENT_BINARY_DIR}/${CONFIG_NAME})
+endforeach()
+
+add_lit_testsuite(check-csan "Running ConcurrencySanitizer tests"
+ ${CSAN_TESTSUITES}
+ DEPENDS ${CSAN_TEST_DEPS})
diff --git a/compiler-rt/test/csan/access-sizes.cpp b/compiler-rt/test/csan/access-sizes.cpp
new file mode 100644
index 00000000000000..9fed52f187b45c
--- /dev/null
+++ b/compiler-rt/test/csan/access-sizes.cpp
@@ -0,0 +1,62 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 1 2>&1 | FileCheck %s --check-prefix=SIZE1
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2 2>&1 | FileCheck %s --check-prefix=SIZE2
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 8 2>&1 | FileCheck %s --check-prefix=SIZE8
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 16 2>&1 | FileCheck %s --check-prefix=SIZE16
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+#include <stdlib.h>
+
+using Vec16 = int __attribute__((vector_size(16)));
+
+#define TEST_SIZE(Name, Type) \
+ volatile Type Name; \
+ static void *Name##Thread(void *) { \
+ Type Value = {}; \
+ RACE_UNTIL_FOUND(I) Name = Value; \
+ return nullptr; \
+ } \
+ static void Name##Test() { \
+ pthread_t T; \
+ pthread_create(&T, nullptr, Name##Thread, nullptr); \
+ Type Value = {}; \
+ RACE_UNTIL_FOUND(I) Name = Value; \
+ pthread_join(T, nullptr); \
+ }
+
+TEST_SIZE(Global1, char)
+TEST_SIZE(Global2, short)
+TEST_SIZE(Global8, long)
+TEST_SIZE(Global16, Vec16)
+
+int main(int Argc, char **Argv) {
+ if (Argc != 2)
+ return 1;
+ switch (atoi(Argv[1])) {
+ case 1:
+ Global1Test();
+ break;
+ case 2:
+ Global2Test();
+ break;
+ case 8:
+ Global8Test();
+ break;
+ case 16:
+ Global16Test();
+ break;
+ default:
+ return 1;
+ }
+ return 0;
+}
+
+// SIZE1: WARNING: ConcurrencySanitizer: data race
+// SIZE1: Write of size 1
+// SIZE2: WARNING: ConcurrencySanitizer: data race
+// SIZE2: Write of size 2
+// SIZE8: WARNING: ConcurrencySanitizer: data race
+// SIZE8: Write of size 8
+// SIZE16: WARNING: ConcurrencySanitizer: data race
+// SIZE16: Write of size 16
diff --git a/compiler-rt/test/csan/atomic-nonatomic.cpp b/compiler-rt/test/csan/atomic-nonatomic.cpp
new file mode 100644
index 00000000000000..8da5a3ccc2b09c
--- /dev/null
+++ b/compiler-rt/test/csan/atomic-nonatomic.cpp
@@ -0,0 +1,27 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+
+int Global;
+
+static void *Thread(void *) {
+ RACE_UNTIL_FOUND(I)
+ __atomic_fetch_add(&Global, 1, __ATOMIC_RELAXED);
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ RACE_UNTIL_FOUND(I)
+ Global++;
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: Write of size 4
+// CHECK: Previous write of size 4
+// CHECK: Thread(void*)
diff --git a/compiler-rt/test/csan/atomic.cpp b/compiler-rt/test/csan/atomic.cpp
new file mode 100644
index 00000000000000..44000912129a08
--- /dev/null
+++ b/compiler-rt/test/csan/atomic.cpp
@@ -0,0 +1,22 @@
+// RUN: %clangxx_csan -O1 %s -o %t -pthread && %run %t 2>&1 | FileCheck %s --allow-empty
+
+#include <pthread.h>
+
+int Global;
+
+static void *Thread(void *) {
+ for (int I = 0; I < 1024; ++I)
+ __atomic_fetch_add(&Global, 1, __ATOMIC_RELAXED);
+ return nullptr;
+}
+
+int main() {
+ pthread_t t;
+ pthread_create(&t, nullptr, Thread, nullptr);
+ for (int I = 0; I < 1024; ++I)
+ __atomic_fetch_add(&Global, 1, __ATOMIC_RELAXED);
+ pthread_join(t, nullptr);
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/halt-on-error.cpp b/compiler-rt/test/csan/halt-on-error.cpp
new file mode 100644
index 00000000000000..c3c15beee5be3a
--- /dev/null
+++ b/compiler-rt/test/csan/halt-on-error.cpp
@@ -0,0 +1,24 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread
+// RUN: not env CSAN_OPTIONS=skip_watch=0:udelay=1000:halt_on_error=1 %run %t 2>&1 | FileCheck %s
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+
+int Global;
+
+static void *Thread(void *) {
+ RACE_UNTIL_FOUND(I)
+ Global++;
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ RACE_UNTIL_FOUND(I)
+ Global++;
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
diff --git a/compiler-rt/test/csan/ignore-thread.cpp b/compiler-rt/test/csan/ignore-thread.cpp
new file mode 100644
index 00000000000000..ca920c655e2459
--- /dev/null
+++ b/compiler-rt/test/csan/ignore-thread.cpp
@@ -0,0 +1,45 @@
+// RUN: %clangxx_csan -O1 %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s
+
+#include <atomic>
+#include <pthread.h>
+#include <stdio.h>
+
+extern "C" void __csan_ignore_thread_begin();
+extern "C" void __csan_ignore_thread_end();
+extern "C" unsigned long long __csan_get_num_data_races();
+
+volatile int Global;
+std::atomic<bool> Start;
+std::atomic<bool> Stop;
+
+__attribute__((no_sanitize("concurrency"))) static void *Thread(void *) {
+ Start.store(true, std::memory_order_release);
+ while (!Stop.load(std::memory_order_relaxed))
+ ++Global;
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ while (!Start.load(std::memory_order_acquire)) {
+ }
+
+ __csan_ignore_thread_begin();
+ __csan_ignore_thread_begin();
+ for (int I = 0; I < 1024; ++I)
+ (void)Global;
+ __csan_ignore_thread_end();
+ for (int I = 0; I < 1024; ++I)
+ (void)Global;
+ __csan_ignore_thread_end();
+
+ Stop.store(true, std::memory_order_relaxed);
+ pthread_join(T, nullptr);
+ printf("races: %llu\n", __csan_get_num_data_races());
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
+// CHECK: races: 0
diff --git a/compiler-rt/test/csan/large-access.cpp b/compiler-rt/test/csan/large-access.cpp
new file mode 100644
index 00000000000000..309e44d9de7b50
--- /dev/null
+++ b/compiler-rt/test/csan/large-access.cpp
@@ -0,0 +1,19 @@
+// RUN: %clangxx_csan -O2 %s -o %t && %run %t 2>&1 | FileCheck %s --allow-empty
+
+using Vec16 = int __attribute__((vector_size(16)));
+using Vec32 = int __attribute__((vector_size(32)));
+
+volatile Vec16 Global16;
+volatile Vec32 Global32;
+
+int main() {
+ Vec16 V16 = {};
+ Vec32 V32 = {};
+ for (int I = 0; I < 1024; ++I) {
+ Global16 = V16;
+ Global32 = V32;
+ }
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/large-range-race.cpp b/compiler-rt/test/csan/large-range-race.cpp
new file mode 100644
index 00000000000000..cefa1cddf95aca
--- /dev/null
+++ b/compiler-rt/test/csan/large-range-race.cpp
@@ -0,0 +1,25 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread && %run %t 2>&1 | FileCheck %s
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+
+char Global[32768];
+
+static void *Thread(void *) {
+ RACE_UNTIL_FOUND(I)
+ __builtin_memset(Global, I, sizeof(Global));
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ RACE_UNTIL_FOUND(I)
+ __builtin_memset(Global, I, sizeof(Global));
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: write of size 8192 at
+// CHECK: Location is global 'Global'
diff --git a/compiler-rt/test/csan/lit.common.cfg.py b/compiler-rt/test/csan/lit.common.cfg.py
new file mode 100644
index 00000000000000..6cc99cadc499ba
--- /dev/null
+++ b/compiler-rt/test/csan/lit.common.cfg.py
@@ -0,0 +1,62 @@
+# -*- Python -*-
+
+import os
+
+
+def get_required_attr(config, attr_name):
+ attr_value = getattr(config, attr_name, None)
+ if attr_value is None:
+ lit_config.fatal("No attribute %r in test configuration!" % attr_name)
+ return attr_value
+
+
+config.name = "CSan-" + config.name_suffix
+config.test_source_root = os.path.dirname(__file__)
+config.suffixes = [".c", ".cpp"]
+
+
+def build_invocation(compile_flags):
+ return " " + " ".join([config.clang] + compile_flags) + " "
+
+
+target_cflags = [get_required_attr(config, "target_cflags")]
+config.substitutions.append(("%clang ", build_invocation(target_cflags)))
+config.substitutions.append(
+ ("%clangxx ", build_invocation(config.cxx_mode_flags + target_cflags))
+)
+config.substitutions.append(
+ ("%clang_csan ", build_invocation(target_cflags + ["-fsanitize=concurrency"]))
+)
+config.substitutions.append(
+ (
+ "%clangxx_csan ",
+ build_invocation(
+ config.cxx_mode_flags + target_cflags + ["-fsanitize=concurrency"]
+ ),
+ )
+)
+
+if config.target_os not in ["Linux"]:
+ config.unsupported = True
+
+if "csan" in config.gpu_runtimes:
+ if "hip" in config.available_features:
+ config.available_features.add("csan-hip")
+ if "openmp-offload" in config.available_features:
+ config.available_features.add("csan-openmp-offload")
+
+
+def add_csan_substitution(name, base):
+ for pattern, replacement in config.substitutions:
+ if pattern == base:
+ config.substitutions.append(
+ (name, replacement.rstrip() + " -fsanitize=concurrency ")
+ )
+ return
+ lit_config.fatal("Missing substitution %r" % base)
+
+
+if "csan-hip" in config.available_features:
+ add_csan_substitution("%clang_hip_csan ", "%clang_hip ")
+if "csan-openmp-offload" in config.available_features:
+ add_csan_substitution("%clang_omp_offload_csan ", "%clang_omp_offload ")
diff --git a/compiler-rt/test/csan/lit.site.cfg.py.in b/compiler-rt/test/csan/lit.site.cfg.py.in
new file mode 100644
index 00000000000000..5fa31b7cc41dbf
--- /dev/null
+++ b/compiler-rt/test/csan/lit.site.cfg.py.in
@@ -0,0 +1,8 @@
+ at LIT_SITE_CFG_IN_HEADER@
+
+config.name_suffix = "@CONFIG_NAME@"
+config.target_cflags = "@CSAN_TEST_TARGET_CFLAGS@"
+config.target_arch = "@CSAN_TEST_TARGET_ARCH@"
+
+lit_config.load_config(config, "@COMPILER_RT_BINARY_DIR@/test/lit.common.configured")
+lit_config.load_config(config, "@CSAN_LIT_TESTS_DIR@/lit.common.cfg.py")
diff --git a/compiler-rt/test/csan/race.cpp b/compiler-rt/test/csan/race.cpp
new file mode 100644
index 00000000000000..b6f060ed2cdd93
--- /dev/null
+++ b/compiler-rt/test/csan/race.cpp
@@ -0,0 +1,26 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread && %run %t 2>&1 | FileCheck %s
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+
+int Global;
+
+static void *Thread(void *) {
+ RACE_UNTIL_FOUND(i)
+ Global++;
+ return nullptr;
+}
+
+int main() {
+ pthread_t t;
+ pthread_create(&t, nullptr, Thread, nullptr);
+ RACE_UNTIL_FOUND(i)
+ Global++;
+ pthread_join(t, nullptr);
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: of size 4 at {{.*}} by thread {{[0-9]+}}:
+// CHECK: #{{[0-9]+}} {{.*}} {{main|Thread}}
+// CHECK: Location is global 'Global'
diff --git a/compiler-rt/test/csan/read-read.cpp b/compiler-rt/test/csan/read-read.cpp
new file mode 100644
index 00000000000000..e3d6dba1a14d5a
--- /dev/null
+++ b/compiler-rt/test/csan/read-read.cpp
@@ -0,0 +1,23 @@
+// RUN: %clangxx_csan -O1 %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=0 %run %t 2>&1 | FileCheck %s --allow-empty
+
+#include <pthread.h>
+
+volatile int Global;
+
+static void *Thread(void *) {
+ for (int I = 0; I < 1 << 16; ++I)
+ (void)Global;
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ for (int I = 0; I < 1 << 16; ++I)
+ (void)Global;
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/single-thread.cpp b/compiler-rt/test/csan/single-thread.cpp
new file mode 100644
index 00000000000000..a01f17e10308ad
--- /dev/null
+++ b/compiler-rt/test/csan/single-thread.cpp
@@ -0,0 +1,11 @@
+// RUN: %clangxx_csan -O1 %s -o %t && %run %t 2>&1 | FileCheck %s --allow-empty
+
+int Global;
+
+int main() {
+ for (int I = 0; I < 1024; ++I)
+ Global++;
+ return 0;
+}
+
+// CHECK-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/unknown-origin.cpp b/compiler-rt/test/csan/unknown-origin.cpp
new file mode 100644
index 00000000000000..383aeff700ab72
--- /dev/null
+++ b/compiler-rt/test/csan/unknown-origin.cpp
@@ -0,0 +1,31 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s
+
+#include <atomic>
+#include <pthread.h>
+
+volatile int Global;
+std::atomic<bool> Start;
+std::atomic<bool> Stop;
+
+__attribute__((no_sanitize("concurrency"))) static void *Thread(void *) {
+ while (!Start.load(std::memory_order_acquire)) {
+ }
+ while (!Stop.load(std::memory_order_relaxed))
+ ++Global;
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ Start.store(true, std::memory_order_release);
+ for (int I = 0; I < 64; ++I)
+ (void)Global;
+ Stop.store(true, std::memory_order_relaxed);
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// CHECK: WARNING: ConcurrencySanitizer: data race of unknown origin
+// CHECK: race at unknown origin, with read of size 4
diff --git a/compiler-rt/test/lit.common.cfg.py b/compiler-rt/test/lit.common.cfg.py
index 64e6ede9b19b40..6ef1517ffe09ef 100644
--- a/compiler-rt/test/lit.common.cfg.py
+++ b/compiler-rt/test/lit.common.cfg.py
@@ -377,6 +377,7 @@ def get_ios_commands_dir():
"MSAN_SYMBOLIZER_PATH",
"LSAN_SYMBOLIZER_PATH",
"UBSAN_SYMBOLIZER_PATH",
+ "CSAN_SYMBOLIZER_PATH",
]
if config.have_disable_symbolizer_path_search:
>From b4a8e89f93697d1f04fc630e39b4ea01fc3cba0b Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 24 Sep 2026 09:56:27 -0500
Subject: [PATCH 2/5] comments
---
.../cmake/Modules/AllSupportedArchDefs.cmake | 1 +
compiler-rt/cmake/config-ix.cmake | 6 ++---
compiler-rt/lib/csan/CMakeLists.txt | 4 ++--
compiler-rt/lib/csan/csan_gpu.cpp | 22 +++++++++++--------
compiler-rt/lib/csan/offload/CMakeLists.txt | 2 +-
compiler-rt/test/csan/CMakeLists.txt | 2 +-
6 files changed, 21 insertions(+), 16 deletions(-)
diff --git a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake
index 9ea6d294a3fe7c..dfb6eafff58c61 100644
--- a/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake
+++ b/compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake
@@ -104,6 +104,7 @@ else()
set(ALL_TSAN_SUPPORTED_ARCH ${X86_64} ${MIPS64} ${ARM64} ${PPC64} ${S390X}
${LOONGARCH64} ${RISCV64})
endif()
+set(ALL_CSAN_SUPPORTED_ARCH ${X86_64} ${AMDGPU})
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}
diff --git a/compiler-rt/cmake/config-ix.cmake b/compiler-rt/cmake/config-ix.cmake
index d1732e761e119b..3167e3c9d36e1f 100644
--- a/compiler-rt/cmake/config-ix.cmake
+++ b/compiler-rt/cmake/config-ix.cmake
@@ -727,6 +727,7 @@ else()
filter_available_targets(PROFILE_SUPPORTED_ARCH ${ALL_PROFILE_SUPPORTED_ARCH})
filter_available_targets(CTX_PROFILE_SUPPORTED_ARCH ${ALL_CTX_PROFILE_SUPPORTED_ARCH})
filter_available_targets(TSAN_SUPPORTED_ARCH ${ALL_TSAN_SUPPORTED_ARCH})
+ filter_available_targets(CSAN_SUPPORTED_ARCH ${ALL_CSAN_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(SAFESTACK_SUPPORTED_ARCH
@@ -897,10 +898,9 @@ else()
set(COMPILER_RT_HAS_UBSAN FALSE)
endif()
-if ((COMPILER_RT_HAS_SANITIZER_COMMON AND UBSAN_SUPPORTED_ARCH AND
+if ((COMPILER_RT_HAS_SANITIZER_COMMON AND CSAN_SUPPORTED_ARCH AND
OS_NAME MATCHES "Linux" AND NOT ANDROID)
- OR (COMPILER_RT_GPU_BUILD AND COMPILER_RT_TARGET_AMDGPU AND
- UBSAN_SUPPORTED_ARCH))
+ OR (COMPILER_RT_GPU_BUILD AND CSAN_SUPPORTED_ARCH))
set(COMPILER_RT_HAS_CSAN TRUE)
else()
set(COMPILER_RT_HAS_CSAN FALSE)
diff --git a/compiler-rt/lib/csan/CMakeLists.txt b/compiler-rt/lib/csan/CMakeLists.txt
index b0791c39e2a4cf..5351ff8b2f2dc2 100644
--- a/compiler-rt/lib/csan/CMakeLists.txt
+++ b/compiler-rt/lib/csan/CMakeLists.txt
@@ -25,7 +25,7 @@ if(COMPILER_RT_GPU_BUILD)
add_compiler_rt_runtime(clang_rt.csan
STATIC
- ARCHS ${UBSAN_SUPPORTED_ARCH}
+ ARCHS ${CSAN_SUPPORTED_ARCH}
SOURCES ${CSAN_SOURCES}
ADDITIONAL_HEADERS ${CSAN_HEADERS}
CFLAGS ${CSAN_CFLAGS}
@@ -42,7 +42,7 @@ add_compiler_rt_component(csan)
add_compiler_rt_runtime(clang_rt.csan
STATIC
- ARCHS ${UBSAN_SUPPORTED_ARCH}
+ ARCHS ${CSAN_SUPPORTED_ARCH}
SOURCES csan.cpp csan_report.cpp
ADDITIONAL_HEADERS csan.h csan_defs.h csan_flags.inc csan_watch.h
OBJECT_LIBS RTSanitizerCommon
diff --git a/compiler-rt/lib/csan/csan_gpu.cpp b/compiler-rt/lib/csan/csan_gpu.cpp
index 729113f0c7950e..5da4d6ddb978ff 100644
--- a/compiler-rt/lib/csan/csan_gpu.cpp
+++ b/compiler-rt/lib/csan/csan_gpu.cpp
@@ -36,7 +36,6 @@ static_assert((CSAN_WATCHPOINT_TABLE_ENTRIES &
static_assert(WP_CHANCE >= 2 && (WP_CHANCE & (WP_CHANCE - 1)) == 0,
"WP_CHANCE must be a power of two");
-// The GPU case does
static constexpr u32 GPU_MAX_ACCESS_SIZE = 16;
static constexpr u32 GPU_WATCHPOINT_ENTRIES = CSAN_WATCHPOINT_TABLE_ENTRIES;
static constexpr u32 GPU_CHECK_ADJACENT_SLOTS = 0;
@@ -70,11 +69,11 @@ INTERFACE u64 __csan_get_num_data_races() {
return __atomic_load_n(&__csan_num_data_races, __ATOMIC_RELAXED);
}
-// Shallow deduplication check to save the host thread work. Keyed on both the
-// PC and the race kind so each distinct kind of race at a PC is reported once.
-static bool should_report(void *pc, unsigned kind) {
+// Shallow deduplication check to save the host thread work. Keyed on the PC
+// pair and the race kind so each distinct race is reported once.
+static bool should_report(uptr pc, uptr peer_pc, unsigned kind) {
static u64 seen[64] = {};
- const u64 token = (reinterpret_cast<uptr>(pc) >> 4) ^
+ const u64 token = (pc >> 4) ^ ((peer_pc >> 4) * 0xD1B54A32D192ED03ull) ^
(static_cast<u64>(kind) * 0x9E3779B97F4A7C15ull);
u64 idx = (token * 0x9E3779B97F4A7C15ull) >> 58;
u64 last = __scoped_atomic_exchange_n(&seen[idx], token, __ATOMIC_RELAXED,
@@ -87,7 +86,7 @@ report(unsigned kind, uptr addr, u32 size, int access_type, uptr pc,
void *peer = nullptr, int peer_access = 0, u32 peer_size = 0,
u8 peer_lane = 0) {
pc = pc ? pc : GET_CALLER_PC();
- if (!should_report(reinterpret_cast<void *>(pc), kind))
+ if (!should_report(pc, reinterpret_cast<uptr>(peer), kind))
return;
__csan_gpu_race rep = {};
@@ -132,7 +131,12 @@ static constexpr u64 TICKS_PER_SEC = 1000000000UL;
// FIXME: Avoids emitting an unresolved reference to the OCLC ABI version.
static u32 num_blocks(int dim) {
#ifdef __AMDGPU__
- return ((const u32 __gpu_constant *)__builtin_amdgcn_implicitarg_ptr())[dim];
+ // The block count excludes a trailing partial work-group.
+ const u8 __gpu_constant *args =
+ (const u8 __gpu_constant *)__builtin_amdgcn_implicitarg_ptr();
+ const u32 count = ((const u32 __gpu_constant *)args)[dim];
+ const u16 remainder = ((const u16 __gpu_constant *)(args + 18))[dim];
+ return count + (remainder > 0);
#else
return __gpu_num_blocks(dim);
#endif
@@ -224,7 +228,7 @@ static u64 read_range(BytePtr bytes, WordPtr, u32 size) {
for (; i < size && ((reinterpret_cast<uptr>(bytes) + i) & 7u); ++i)
sum = (sum ^ bytes[i]) * 0x100000001b3ull;
- for (; i + 8 <= size; i += 8)
+ for (; size - i >= 8; i += 8)
sum = (sum ^ *reinterpret_cast<WordPtr>(bytes + i)) * 0x100000001b3ull;
for (; i < size; ++i)
sum = (sum ^ bytes[i]) * 0x100000001b3ull;
@@ -410,7 +414,7 @@ static void check_access(const volatile void *addr, uptr size, int access_type,
// Public ABI (emitted by the ConcurrencySanitizer pass)
//===----------------------------------------------------------------------===//
-// Using `sanitize_concurrency_no_checking_at_run_time` ignored on the device,
+// The device does not support `sanitize_concurrency_no_checking_at_run_time`.
INTERFACE void __csan_init() {}
INTERFACE void __csan_func_entry(void *) {}
INTERFACE void __csan_func_exit() {}
diff --git a/compiler-rt/lib/csan/offload/CMakeLists.txt b/compiler-rt/lib/csan/offload/CMakeLists.txt
index 346be95f20fb96..338177e1ac74e3 100644
--- a/compiler-rt/lib/csan/offload/CMakeLists.txt
+++ b/compiler-rt/lib/csan/offload/CMakeLists.txt
@@ -16,7 +16,7 @@ set(CSAN_OFFLOAD_HEADERS
add_compiler_rt_runtime(clang_rt.csan_offload
STATIC
- ARCHS ${UBSAN_SUPPORTED_ARCH}
+ ARCHS ${CSAN_SUPPORTED_ARCH}
SOURCES ${CSAN_OFFLOAD_SOURCES}
ADDITIONAL_HEADERS ${CSAN_OFFLOAD_HEADERS}
OBJECT_LIBS RTSanitizerOffload
diff --git a/compiler-rt/test/csan/CMakeLists.txt b/compiler-rt/test/csan/CMakeLists.txt
index 6df0f1f1de3d11..c8bd5d3e06e1e1 100644
--- a/compiler-rt/test/csan/CMakeLists.txt
+++ b/compiler-rt/test/csan/CMakeLists.txt
@@ -9,7 +9,7 @@ set(CSAN_TESTSUITES)
set(CSAN_TEST_DEPS ${SANITIZER_COMMON_LIT_TEST_DEPS})
list(APPEND CSAN_TEST_DEPS csan)
-set(CSAN_TEST_ARCH ${UBSAN_SUPPORTED_ARCH})
+set(CSAN_TEST_ARCH ${CSAN_SUPPORTED_ARCH})
foreach(arch ${CSAN_TEST_ARCH})
set(CSAN_TEST_TARGET_ARCH ${arch})
get_test_cc_for_arch(${arch} CSAN_TEST_TARGET_CC CSAN_TEST_TARGET_CFLAGS)
>From a49fa31afef678b9bd9c24c515757f8a477d80dc Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 29 Sep 2026 13:51:55 -0500
Subject: [PATCH 3/5] [compiler-rt] Remove dlsym interceptor and support
`-shared-libsan` for CSan
Summary:
Follow the UBSan offload runtime. Offload now resolves HSA through the
global scope, so the `dlsym` interceptor is no longer needed. The real
HSA entry points are still taken from the loaded HSA library rather than
`RTLD_NEXT`, since every DSO with a static runtime exports the same
wrappers and they would otherwise chain back into each other.
Build `libclang_rt.csan.so` with the offload objects folded in. The
exported HSA wrappers report failure when HSA is absent and warn when HSA
was loaded ahead of the runtime. The preinit hook moves to a separate
`csan_offload-preinit` archive for executables.
---
compiler-rt/lib/csan/CMakeLists.txt | 93 ++++++++++--
compiler-rt/lib/csan/csan.syms.extra | 1 +
compiler-rt/lib/csan/offload/CMakeLists.txt | 24 +++-
.../offload/csan_offload_hsa_interceptors.cpp | 132 +++++-------------
.../lib/csan/offload/csan_offload_preinit.cpp | 22 +++
compiler-rt/test/csan/AMDGPU/global-race.hip | 9 ++
compiler-rt/test/csan/AMDGPU/openmp-race.cpp | 2 +
compiler-rt/test/csan/hsa-load-order.cpp | 35 +++++
compiler-rt/test/csan/hsa-missing.cpp | 24 ++++
compiler-rt/test/csan/lit.common.cfg.py | 6 +
compiler-rt/test/csan/race.cpp | 2 +
11 files changed, 240 insertions(+), 110 deletions(-)
create mode 100644 compiler-rt/lib/csan/csan.syms.extra
create mode 100644 compiler-rt/lib/csan/offload/csan_offload_preinit.cpp
create mode 100644 compiler-rt/test/csan/hsa-load-order.cpp
create mode 100644 compiler-rt/test/csan/hsa-missing.cpp
diff --git a/compiler-rt/lib/csan/CMakeLists.txt b/compiler-rt/lib/csan/CMakeLists.txt
index 5351ff8b2f2dc2..4ae3401f8c2dd9 100644
--- a/compiler-rt/lib/csan/CMakeLists.txt
+++ b/compiler-rt/lib/csan/CMakeLists.txt
@@ -35,26 +35,101 @@ if(COMPILER_RT_GPU_BUILD)
return()
endif()
+set(CSAN_SOURCES
+ csan.cpp
+ csan_report.cpp
+ )
+
+set(CSAN_HEADERS
+ csan.h
+ csan_defs.h
+ csan_flags.inc
+ csan_watch.h
+ )
+
set(CSAN_CFLAGS ${SANITIZER_COMMON_CFLAGS})
append_rtti_flag(OFF CSAN_CFLAGS)
+set(CSAN_LINK_FLAGS ${SANITIZER_COMMON_LINK_FLAGS})
+
+set(CSAN_DYNAMIC_LIBS
+ ${COMPILER_RT_UNWINDER_LINK_LIBS}
+ ${SANITIZER_CXX_ABI_LIBRARIES}
+ ${SANITIZER_COMMON_LINK_LIBS})
+
+append_list_if(COMPILER_RT_HAS_LIBDL dl CSAN_DYNAMIC_LIBS)
+append_list_if(COMPILER_RT_HAS_LIBRT rt CSAN_DYNAMIC_LIBS)
+append_list_if(COMPILER_RT_HAS_LIBPTHREAD pthread CSAN_DYNAMIC_LIBS)
+if (COMPILER_RT_ENABLE_INTERNAL_SYMBOLIZER)
+ append_list_if(COMPILER_RT_HAS_LIBM m CSAN_DYNAMIC_LIBS)
+endif()
+
+set(CSAN_COMMON_RUNTIME_OBJECT_LIBS
+ RTSanitizerCommon
+ RTSanitizerCommonLibc
+ RTSanitizerCommonCoverage
+ RTSanitizerCommonSymbolizer
+ RTSanitizerCommonSymbolizerInternal
+ RTInterception)
+
add_compiler_rt_component(csan)
+add_compiler_rt_object_libraries(RTCsan
+ ARCHS ${CSAN_SUPPORTED_ARCH}
+ SOURCES ${CSAN_SOURCES}
+ ADDITIONAL_HEADERS ${CSAN_HEADERS}
+ CFLAGS ${CSAN_CFLAGS})
+
add_compiler_rt_runtime(clang_rt.csan
STATIC
ARCHS ${CSAN_SUPPORTED_ARCH}
- SOURCES csan.cpp csan_report.cpp
- ADDITIONAL_HEADERS csan.h csan_defs.h csan_flags.inc csan_watch.h
- OBJECT_LIBS RTSanitizerCommon
- RTSanitizerCommonLibc
- RTSanitizerCommonCoverage
- RTSanitizerCommonSymbolizer
- RTSanitizerCommonSymbolizerInternal
- RTInterception
+ OBJECT_LIBS RTCsan
+ ${CSAN_COMMON_RUNTIME_OBJECT_LIBS}
CFLAGS ${CSAN_CFLAGS}
PARENT_TARGET csan)
-# Thin host interceptors for offloading. Linked with -u to keep it separate.
+# 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_VERSION_SCRIPT)
+ file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/dummy.cpp "")
+ add_compiler_rt_object_libraries(RTCsan_dynamic_version_script_dummy
+ ARCHS ${CSAN_SUPPORTED_ARCH}
+ SOURCES ${CMAKE_CURRENT_BINARY_DIR}/dummy.cpp
+ CFLAGS ${CSAN_CFLAGS})
+
+ foreach(arch ${CSAN_SUPPORTED_ARCH})
+ set(CSAN_OFFLOAD_OBJECT_LIBS)
+ set(CSAN_OFFLOAD_VERSION_LIBS)
+ if(TARGET RTCsan_offload.${arch} AND TARGET RTSanitizerOffload.${arch})
+ set(CSAN_OFFLOAD_OBJECT_LIBS RTCsan_offload RTSanitizerOffload)
+ set(CSAN_OFFLOAD_VERSION_LIBS clang_rt.csan_offload-${arch})
+ endif()
+
+ add_sanitizer_rt_version_list(clang_rt.csan-dynamic-${arch}
+ LIBS clang_rt.csan-${arch}
+ ${CSAN_OFFLOAD_VERSION_LIBS}
+ EXTRA csan.syms.extra)
+ set(VERSION_SCRIPT_FLAG
+ -Wl,--version-script,${CMAKE_CURRENT_BINARY_DIR}/clang_rt.csan-dynamic-${arch}.vers)
+ set_property(SOURCE
+ ${CMAKE_CURRENT_BINARY_DIR}/dummy.cpp
+ APPEND PROPERTY
+ OBJECT_DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/clang_rt.csan-dynamic-${arch}.vers)
+
+ add_compiler_rt_runtime(clang_rt.csan
+ SHARED
+ ARCHS ${arch}
+ OBJECT_LIBS RTCsan
+ ${CSAN_COMMON_RUNTIME_OBJECT_LIBS}
+ ${CSAN_OFFLOAD_OBJECT_LIBS}
+ RTCsan_dynamic_version_script_dummy
+ CFLAGS ${CSAN_CFLAGS}
+ LINK_FLAGS ${CSAN_LINK_FLAGS} ${VERSION_SCRIPT_FLAG}
+ LINK_LIBS ${CSAN_DYNAMIC_LIBS}
+ PARENT_TARGET csan)
+ endforeach()
+endif()
diff --git a/compiler-rt/lib/csan/csan.syms.extra b/compiler-rt/lib/csan/csan.syms.extra
new file mode 100644
index 00000000000000..b60f39ee0742f3
--- /dev/null
+++ b/compiler-rt/lib/csan/csan.syms.extra
@@ -0,0 +1 @@
+__csan_*
diff --git a/compiler-rt/lib/csan/offload/CMakeLists.txt b/compiler-rt/lib/csan/offload/CMakeLists.txt
index 338177e1ac74e3..76330eae549cf4 100644
--- a/compiler-rt/lib/csan/offload/CMakeLists.txt
+++ b/compiler-rt/lib/csan/offload/CMakeLists.txt
@@ -14,12 +14,30 @@ set(CSAN_OFFLOAD_HEADERS
../csan_offload_packet.h
)
+add_compiler_rt_object_libraries(RTCsan_offload
+ ARCHS ${CSAN_SUPPORTED_ARCH}
+ SOURCES ${CSAN_OFFLOAD_SOURCES}
+ ADDITIONAL_HEADERS ${CSAN_OFFLOAD_HEADERS}
+ CFLAGS ${CSAN_CFLAGS})
+foreach(arch ${CSAN_SUPPORTED_ARCH})
+ target_link_libraries(RTCsan_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.csan_offload
STATIC
ARCHS ${CSAN_SUPPORTED_ARCH}
- SOURCES ${CSAN_OFFLOAD_SOURCES}
+ OBJECT_LIBS RTCsan_offload
+ RTSanitizerOffload
+ CFLAGS ${CSAN_CFLAGS}
+ PARENT_TARGET csan)
+
+add_compiler_rt_runtime(clang_rt.csan_offload-preinit
+ STATIC
+ ARCHS ${CSAN_SUPPORTED_ARCH}
+ SOURCES csan_offload_preinit.cpp
ADDITIONAL_HEADERS ${CSAN_OFFLOAD_HEADERS}
- OBJECT_LIBS RTSanitizerOffload
CFLAGS ${CSAN_CFLAGS}
- LINK_LIBS llvm-libc-common-utilities
PARENT_TARGET csan)
diff --git a/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp b/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
index c08dc58a49637d..7d8b4ca921340a 100644
--- a/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
+++ b/compiler-rt/lib/csan/offload/csan_offload_hsa_interceptors.cpp
@@ -30,10 +30,6 @@
#error "Offload CSan reporting is supported on Linux only"
#endif
-#if SANITIZER_GLIBC
-#pragma weak dlvsym
-#endif
-
using namespace __sanitizer;
using namespace __csan;
@@ -77,25 +73,39 @@ void Initialize() {
Offload::Get().RegisterHandler(HandleOffloadReport);
Atexit([] { Offload::Get().UntrackImages(); });
AddDieCallback([] { Offload::Get().UntrackImages(); });
-
- // Mark ready before LateInitialize, as it can be reentrant through dlsym.
atomic_store(&Initialized, 1, memory_order_release);
Symbolizer::LateInitialize();
}
} // namespace
-static void BindRealDlsym();
-static void *HsaSymbol(const char *Name);
+// DSOs with a static runtime all export these wrappers and interpose onto the
+// first, so 'RTLD_NEXT' can loop back into it. Resolve from HSA directly
+// without loading it.
+static void *HsaSymbol(const char *Name) {
+ SpinMutexLock L(&HsaMutex);
+ constexpr const char *Libs[] = {SANITIZER_HSA_LIBRARY ".so.1",
+ SANITIZER_HSA_LIBRARY ".so"};
+ for (const char *Lib : Libs)
+ if (!HsaHandle)
+ HsaHandle = dlopen(Lib, RTLD_LAZY | RTLD_NOLOAD);
+ return HsaHandle ? dlsym(HsaHandle, Name) : nullptr;
+}
+
+template <typename T> static T HsaFunction(const char *Name) {
+ return reinterpret_cast<T>(HsaSymbol(Name));
+}
+// The shared runtime exports these to every program, act as if HSA is absent
+// when it is not loaded.
#define CSAN_HSA_ENTER(name) \
Initialize(); \
- if (UNLIKELY(!REAL(name) || REAL(name) == name)) { \
- REAL(name) = reinterpret_cast<decltype(REAL(name))>(HsaSymbol(#name)); \
- if (UNLIKELY(!REAL(name) || REAL(name) == name)) { \
- Report("ERROR: %s: cannot find %s in this process\n", SanitizerToolName, \
- #name); \
- Die(); \
+ if (UNLIKELY(!REAL(name))) { \
+ REAL(name) = HsaFunction<decltype(REAL(name))>(#name); \
+ if (UNLIKELY(!REAL(name))) { \
+ VReport(1, "%s: cannot find %s in this process\n", SanitizerToolName, \
+ #name); \
+ return HSA_STATUS_ERROR; \
} \
}
@@ -104,23 +114,6 @@ static void *HsaSymbol(const char *Name);
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
-#define CSAN_HSA_WRAPS(X) \
- X(hsa_init) \
- X(hsa_shut_down) \
- X(hsa_executable_freeze) \
- X(hsa_executable_destroy)
-
-static void *WrapperFor(const char *Name) {
-#define CSAN_HSA_WRAP(Fn) \
- if (!internal_strcmp(Name, #Fn)) \
- return reinterpret_cast<void *>(Fn);
- CSAN_HSA_WRAPS(CSAN_HSA_WRAP)
-#undef CSAN_HSA_WRAP
- return nullptr;
-}
-
static bool FromHsa(void *P) {
Dl_info Info = {};
if (!dladdr(P, &Info) || !Info.dli_fname)
@@ -128,69 +121,16 @@ static bool FromHsa(void *P) {
return internal_strstr(Info.dli_fname, SANITIZER_HSA_LIBRARY);
}
-// 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]]
- return REAL(dlsym)(Handle, Name);
-
- void *Sym = REAL(dlsym)(Handle, Name);
- if (!Sym || !Name)
- return Sym;
-
- void *Wrapper = WrapperFor(Name);
- if (!Wrapper || !FromHsa(Sym))
- return Sym;
- return Wrapper;
-}
-#else
-DEFINE_REAL(void *, dlsym, void *, const char *)
-#endif
-
-static void BindRealDlsym() {
- if (LIKELY(REAL(dlsym)))
+// Callers bind to whichever 'hsa_init' comes first, if HSA was loaded before
+// the runtime the interceptors are bypassed.
+static void CheckInterposed() {
+ void *Sym = dlsym(RTLD_DEFAULT, "hsa_init");
+ if (!Sym || !FromHsa(Sym))
return;
-#if SANITIZER_GLIBC
- static const char *kVers[] = {"GLIBC_2.34", "GLIBC_2.17", "GLIBC_2.2.5",
- "GLIBC_2.0"};
- if (dlvsym) {
- for (const char *Ver : kVers) {
- if (void *P = dlvsym(RTLD_NEXT, "dlsym", Ver)) {
- REAL(dlsym) = reinterpret_cast<decltype(REAL(dlsym))>(P);
- return;
- }
- }
- }
-#endif
- Report("ERROR: %s: cannot bind dlsym\n", SanitizerToolName);
- Die();
-}
-
-static void *HsaSymbol(const char *Name) {
- BindRealDlsym();
- if (!HsaHandle) {
- SpinMutexLock L(&HsaMutex);
- if (HsaHandle)
- return REAL(dlsym)(HsaHandle, Name);
- constexpr const char *Names[] = {"libhsa-runtime64.so.1",
- "libhsa-runtime64.so"};
- for (const char *Name : Names)
- if (void *H = dlopen(Name, RTLD_LAZY | RTLD_NOLOAD))
- HsaHandle = H;
- for (const char *Name : Names)
- if (!HsaHandle)
- HsaHandle = dlopen(Name, RTLD_LAZY | RTLD_LOCAL);
- }
- return HsaHandle ? REAL(dlsym)(HsaHandle, Name) : nullptr;
-}
-
-template <typename T> static T HsaFunction(const char *Name) {
- return reinterpret_cast<T>(HsaSymbol(Name));
+ Report("WARNING: %s: the runtime is loaded too late to intercept HSA, GPU "
+ "races will not be reported. Link the runtime first or use "
+ "LD_PRELOAD.\n",
+ SanitizerToolName);
}
static bool Lookup(hsa_executable_t Executable, const char *Name,
@@ -360,11 +300,7 @@ INTERCEPTOR(hsa_status_t, hsa_executable_destroy, hsa_executable_t Executable) {
extern "C" void __csan_offload_init() { Initialize(); }
-#if SANITIZER_CAN_USE_PREINIT_ARRAY
-__attribute__((section(".preinit_array"), used)) static void (
- *csan_offload_preinit)(void) = __csan_offload_init;
-#endif
-
__attribute__((constructor(0))) static void CsanOffloadDynInit() {
__csan_offload_init();
+ CheckInterposed();
}
diff --git a/compiler-rt/lib/csan/offload/csan_offload_preinit.cpp b/compiler-rt/lib/csan/offload/csan_offload_preinit.cpp
new file mode 100644
index 00000000000000..c060c9edac55d6
--- /dev/null
+++ b/compiler-rt/lib/csan/offload/csan_offload_preinit.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Call __csan_offload_init at the very early stage of process startup.
+///
+//===----------------------------------------------------------------------===//
+
+#include "csan_offload.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+
+#if SANITIZER_CAN_USE_PREINIT_ARRAY
+// This section is linked into the main executable when offloading uses CSan
+// to perform initialization at a very early stage.
+__attribute__((section(".preinit_array"), used)) static auto preinit =
+ __csan_offload_init;
+#endif
diff --git a/compiler-rt/test/csan/AMDGPU/global-race.hip b/compiler-rt/test/csan/AMDGPU/global-race.hip
index 90970929800f0f..2ba58c58c6aa07 100644
--- a/compiler-rt/test/csan/AMDGPU/global-race.hip
+++ b/compiler-rt/test/csan/AMDGPU/global-race.hip
@@ -2,6 +2,15 @@
// RUN: %run %t 2>&1 | FileCheck %s
// RUN: %clang_hip_csan -O2 %s -o %t %hip_libs
// RUN: %run %t 2>&1 | FileCheck %s
+// RUN: %clang_hip_csan -shared-libsan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s
+
+// The host has no references to the shared runtime, it must still be loaded.
+// RUN: %clang_hip -Xarch_device -fsanitize=concurrency -shared-libsan %s \
+// RUN: -o %t %hip_libs
+// RUN: llvm-readelf -d %t | FileCheck %s --check-prefix=NEEDED
+// RUN: %run %t 2>&1 | FileCheck %s
+// NEEDED: libclang_rt.csan.so
#include "race.h"
diff --git a/compiler-rt/test/csan/AMDGPU/openmp-race.cpp b/compiler-rt/test/csan/AMDGPU/openmp-race.cpp
index 2c053c0e8e6681..b3b4f301d3a2ab 100644
--- a/compiler-rt/test/csan/AMDGPU/openmp-race.cpp
+++ b/compiler-rt/test/csan/AMDGPU/openmp-race.cpp
@@ -1,5 +1,7 @@
// RUN: %clang_omp_offload_csan -fno-exceptions %s -o %t
// RUN: %run %t 2>&1 | FileCheck %s
+// RUN: %clang_omp_offload_csan -fno-exceptions -shared-libsan %s -o %t
+// RUN: %run %t 2>&1 | FileCheck %s
#include "race.h"
diff --git a/compiler-rt/test/csan/hsa-load-order.cpp b/compiler-rt/test/csan/hsa-load-order.cpp
new file mode 100644
index 00000000000000..18ecf6af6f8c0d
--- /dev/null
+++ b/compiler-rt/test/csan/hsa-load-order.cpp
@@ -0,0 +1,35 @@
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: %clangxx -DBUILD_HSA -fPIC -shared %s -o %t/libhsa-runtime64.so.1 \
+// RUN: -Wl,-soname,libhsa-runtime64.so.1
+// RUN: %clangxx_csan -DBUILD_DSO -shared-libsan -fPIC -shared %s \
+// RUN: -o %t/libdso.so %t/libhsa-runtime64.so.1 -Wl,-rpath,%t
+// RUN: %clangxx %s -o %t/early %t/libdso.so -Wl,-rpath,%t
+// RUN: %run %t/early 2>&1 | FileCheck %s --check-prefix=EARLY
+// RUN: %clangxx %s -o %t/late %t/libhsa-runtime64.so.1 %t/libdso.so \
+// RUN: -Wl,-rpath,%t
+// RUN: %run %t/late 2>&1 | FileCheck %s --check-prefix=LATE
+
+// The shared runtime must warn if HSA is loaded ahead of it, since the
+// interceptors are then bypassed.
+
+// REQUIRES: csan-offload
+
+#include <stdio.h>
+
+#if defined(BUILD_HSA)
+extern "C" int hsa_init() { return 42; }
+#elif defined(BUILD_DSO)
+extern "C" int dso() { return 0; }
+#else
+extern "C" int dso();
+int main() {
+ fprintf(stderr, "DONE\n");
+ return dso();
+}
+#endif
+
+// EARLY-NOT: WARNING
+// EARLY: DONE
+
+// LATE: WARNING: ConcurrencySanitizer: the runtime is loaded too late
+// LATE: DONE
diff --git a/compiler-rt/test/csan/hsa-missing.cpp b/compiler-rt/test/csan/hsa-missing.cpp
new file mode 100644
index 00000000000000..5d178807f5437b
--- /dev/null
+++ b/compiler-rt/test/csan/hsa-missing.cpp
@@ -0,0 +1,24 @@
+// RUN: %clangxx_csan -shared-libsan %s -o %t
+// RUN: %run %t 2>&1 | FileCheck %s
+
+// The shared runtime exports the HSA wrappers to every program. Without HSA
+// loaded they must report failure rather than kill the process.
+
+// REQUIRES: csan-offload
+
+#include <dlfcn.h>
+#include <stdio.h>
+
+int main() {
+ auto Init = reinterpret_cast<int (*)()>(dlsym(RTLD_DEFAULT, "hsa_init"));
+ auto ShutDown =
+ reinterpret_cast<int (*)()>(dlsym(RTLD_DEFAULT, "hsa_shut_down"));
+ if (!Init || !ShutDown)
+ return 1;
+ fprintf(stderr, "init: %s\n", Init() ? "failed" : "succeeded");
+ fprintf(stderr, "shut_down: %s\n", ShutDown() ? "failed" : "succeeded");
+ return 0;
+}
+
+// CHECK: init: failed
+// CHECK: shut_down: failed
diff --git a/compiler-rt/test/csan/lit.common.cfg.py b/compiler-rt/test/csan/lit.common.cfg.py
index 6cc99cadc499ba..0c6ba0a75a714a 100644
--- a/compiler-rt/test/csan/lit.common.cfg.py
+++ b/compiler-rt/test/csan/lit.common.cfg.py
@@ -39,6 +39,12 @@ def build_invocation(compile_flags):
if config.target_os not in ["Linux"]:
config.unsupported = True
+# The host offloading runtime is optional, see lib/csan/offload.
+if os.path.exists(
+ os.path.join(config.compiler_rt_libdir, "libclang_rt.csan_offload.a")
+):
+ config.available_features.add("csan-offload")
+
if "csan" in config.gpu_runtimes:
if "hip" in config.available_features:
config.available_features.add("csan-hip")
diff --git a/compiler-rt/test/csan/race.cpp b/compiler-rt/test/csan/race.cpp
index b6f060ed2cdd93..f0dfcd848cca59 100644
--- a/compiler-rt/test/csan/race.cpp
+++ b/compiler-rt/test/csan/race.cpp
@@ -1,4 +1,6 @@
// RUN: %clangxx_csan -O1 -g %s -o %t -pthread && %run %t 2>&1 | FileCheck %s
+// RUN: %clangxx_csan -O1 -g -shared-libsan %s -o %t -pthread
+// RUN: %run %t 2>&1 | FileCheck %s
#include "AMDGPU/race.h"
#include <pthread.h>
>From 20e22f9b933257575714e803e750033de3bd511b Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 1 Oct 2026 17:01:52 -0500
Subject: [PATCH 4/5] Remove separate volatile entry points
---
compiler-rt/lib/csan/csan.cpp | 4 --
compiler-rt/lib/csan/csan_gpu.cpp | 4 --
compiler-rt/test/csan/AMDGPU/volatile.hip | 22 ++++++++++
compiler-rt/test/csan/volatile.cpp | 49 +++++++++++++++++++++++
4 files changed, 71 insertions(+), 8 deletions(-)
create mode 100644 compiler-rt/test/csan/AMDGPU/volatile.hip
create mode 100644 compiler-rt/test/csan/volatile.cpp
diff --git a/compiler-rt/lib/csan/csan.cpp b/compiler-rt/lib/csan/csan.cpp
index 148e30643333ae..58eea09a567848 100644
--- a/compiler-rt/lib/csan/csan.cpp
+++ b/compiler-rt/lib/csan/csan.cpp
@@ -304,12 +304,8 @@ static int AccessFlags(int Flags, bool IsWrite) {
#define CSAN_ACCESS(N) \
CSAN_PROBE(__csan_read##N, N, false) \
CSAN_PROBE(__csan_unaligned_read##N, N, false) \
- CSAN_PROBE(__csan_volatile_read##N, N, false) \
- CSAN_PROBE(__csan_unaligned_volatile_read##N, N, false) \
CSAN_PROBE(__csan_write##N, N, true) \
CSAN_PROBE(__csan_unaligned_write##N, N, true) \
- CSAN_PROBE(__csan_volatile_write##N, N, true) \
- CSAN_PROBE(__csan_unaligned_volatile_write##N, N, true) \
CSAN_PROBE(__csan_read_write##N, N, true) \
CSAN_PROBE(__csan_unaligned_read_write##N, N, true)
diff --git a/compiler-rt/lib/csan/csan_gpu.cpp b/compiler-rt/lib/csan/csan_gpu.cpp
index 5da4d6ddb978ff..cd881dd3362a2a 100644
--- a/compiler-rt/lib/csan/csan_gpu.cpp
+++ b/compiler-rt/lib/csan/csan_gpu.cpp
@@ -433,12 +433,8 @@ static int access_flags(int flags, bool is_write) {
#define CSAN_ACCESS(N) \
CSAN_PROBE(__csan_read##N, N, false) \
CSAN_PROBE(__csan_unaligned_read##N, N, false) \
- CSAN_PROBE(__csan_volatile_read##N, N, false) \
- CSAN_PROBE(__csan_unaligned_volatile_read##N, N, false) \
CSAN_PROBE(__csan_write##N, N, true) \
CSAN_PROBE(__csan_unaligned_write##N, N, true) \
- CSAN_PROBE(__csan_volatile_write##N, N, true) \
- CSAN_PROBE(__csan_unaligned_volatile_write##N, N, true) \
CSAN_PROBE(__csan_read_write##N, N, true) \
CSAN_PROBE(__csan_unaligned_read_write##N, N, true)
diff --git a/compiler-rt/test/csan/AMDGPU/volatile.hip b/compiler-rt/test/csan/AMDGPU/volatile.hip
new file mode 100644
index 00000000000000..2fae11a78bead3
--- /dev/null
+++ b/compiler-rt/test/csan/AMDGPU/volatile.hip
@@ -0,0 +1,22 @@
+// RUN: %clang_hip_csan %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --check-prefix=RACE
+// RUN: %clang_hip_csan -mllvm -csan-distinguish-volatile %s -o %t %hip_libs
+// RUN: %run %t 2>&1 | FileCheck %s --check-prefix=MARKED --allow-empty
+
+#include "race.h"
+
+__device__ volatile int Data;
+
+__global__ void kernel() {
+ for (int I = 0; I < 1 << 14 && !race_found(); ++I)
+ Data = I;
+}
+
+int main() {
+ kernel<<<64, 64>>>();
+ CHECK_HIP(hipDeviceSynchronize());
+ return 0;
+}
+
+// RACE: WARNING: ConcurrencySanitizer: {{.*}}race
+// MARKED-NOT: WARNING: ConcurrencySanitizer
diff --git a/compiler-rt/test/csan/volatile.cpp b/compiler-rt/test/csan/volatile.cpp
new file mode 100644
index 00000000000000..2c3024ed413f99
--- /dev/null
+++ b/compiler-rt/test/csan/volatile.cpp
@@ -0,0 +1,49 @@
+// RUN: %clangxx_csan -O1 -g %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s --check-prefix=RACE
+// RUN: %clangxx_csan -O1 -g -mllvm -csan-distinguish-volatile %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s --check-prefix=MARKED --allow-empty
+// RUN: %clangxx_csan -O1 -g -mllvm -csan-distinguish-volatile -DPLAIN %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s --check-prefix=RACE
+// RUN: %clangxx_csan -O1 -g -mllvm -csan-distinguish-volatile -DUNALIGNED %s -o %t -pthread
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2>&1 | FileCheck %s --check-prefix=RACE
+
+#include "AMDGPU/race.h"
+#include <pthread.h>
+
+struct __attribute__((packed)) Packed {
+ char Pad;
+ volatile int Value;
+};
+
+int Global;
+alignas(8) Packed Unaligned;
+
+#if defined(UNALIGNED)
+# define VOLATILE_STORE(V) (Unaligned.Value = (V))
+#else
+# define VOLATILE_STORE(V) (*(volatile int *)&Global = (V))
+#endif
+
+#if defined(PLAIN)
+# define MAIN_STORE(V) (Global = (V))
+#else
+# define MAIN_STORE(V) VOLATILE_STORE(V)
+#endif
+
+static void *Thread(void *) {
+ RACE_UNTIL_FOUND(I)
+ VOLATILE_STORE(I);
+ return nullptr;
+}
+
+int main() {
+ pthread_t T;
+ pthread_create(&T, nullptr, Thread, nullptr);
+ RACE_UNTIL_FOUND(I)
+ MAIN_STORE(I);
+ pthread_join(T, nullptr);
+ return 0;
+}
+
+// RACE: WARNING: ConcurrencySanitizer: data race
+// MARKED-NOT: WARNING: ConcurrencySanitizer
>From 229beb82ab1417e00db7e69a6da09dcbdeeabc55 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 1 Oct 2026 19:50:46 -0500
Subject: [PATCH 5/5] Use unaligned loads when snapshotting watched values
---
compiler-rt/lib/csan/csan.cpp | 6 +++---
compiler-rt/test/csan/access-sizes.cpp | 22 ++++++++++++++++++++++
2 files changed, 25 insertions(+), 3 deletions(-)
diff --git a/compiler-rt/lib/csan/csan.cpp b/compiler-rt/lib/csan/csan.cpp
index 58eea09a567848..31d6d3032ebe25 100644
--- a/compiler-rt/lib/csan/csan.cpp
+++ b/compiler-rt/lib/csan/csan.cpp
@@ -170,11 +170,11 @@ static u64 ReadInstrumented(const volatile void *Ptr, uptr Size) {
case 1:
return *(const volatile u8 *)Ptr;
case 2:
- return *(const volatile u16 *)Ptr;
+ return *(const volatile uu16 *)Ptr;
case 4:
- return *(const volatile u32 *)Ptr;
+ return *(const volatile uu32 *)Ptr;
case 8:
- return *(const volatile u64 *)Ptr;
+ return *(const volatile uu64 *)Ptr;
default:
return ReadRange((const volatile u8 *)Ptr, Size);
}
diff --git a/compiler-rt/test/csan/access-sizes.cpp b/compiler-rt/test/csan/access-sizes.cpp
index 9fed52f187b45c..993616bce74eeb 100644
--- a/compiler-rt/test/csan/access-sizes.cpp
+++ b/compiler-rt/test/csan/access-sizes.cpp
@@ -3,6 +3,7 @@
// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 2 2>&1 | FileCheck %s --check-prefix=SIZE2
// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 8 2>&1 | FileCheck %s --check-prefix=SIZE8
// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 16 2>&1 | FileCheck %s --check-prefix=SIZE16
+// RUN: env CSAN_OPTIONS=skip_watch=0:udelay=1000 %run %t 0 2>&1 | FileCheck %s --check-prefix=UNALIGNED
#include "AMDGPU/race.h"
#include <pthread.h>
@@ -30,10 +31,29 @@ TEST_SIZE(Global2, short)
TEST_SIZE(Global8, long)
TEST_SIZE(Global16, Vec16)
+struct __attribute__((packed, aligned(4))) Packed {
+ char Pad;
+ int Value;
+};
+volatile Packed Unaligned;
+static void *UnalignedThread(void *) {
+ RACE_UNTIL_FOUND(I) Unaligned.Value = 0;
+ return nullptr;
+}
+static void UnalignedTest() {
+ pthread_t T;
+ pthread_create(&T, nullptr, UnalignedThread, nullptr);
+ RACE_UNTIL_FOUND(I) Unaligned.Value = 0;
+ pthread_join(T, nullptr);
+}
+
int main(int Argc, char **Argv) {
if (Argc != 2)
return 1;
switch (atoi(Argv[1])) {
+ case 0:
+ UnalignedTest();
+ break;
case 1:
Global1Test();
break;
@@ -60,3 +80,5 @@ int main(int Argc, char **Argv) {
// SIZE8: Write of size 8
// SIZE16: WARNING: ConcurrencySanitizer: data race
// SIZE16: Write of size 16
+// UNALIGNED: WARNING: ConcurrencySanitizer: data race
+// UNALIGNED: Write of size 4
More information about the llvm-branch-commits
mailing list