[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
Wed Sep 23 11:55:11 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/225782
>From be40d5b0705b46f10ea3fc91115473f04dbb0aad 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] [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_common/sanitizer_offload.cpp | 49 +-
.../lib/sanitizer_common/sanitizer_offload.h | 3 +
.../sanitizer_common/sanitizer_offload_hsa.h | 2 +
.../sanitizer_offload_image.cpp | 21 +
.../sanitizer_offload_opcodes.h | 1 +
.../sanitizer_common/sanitizer_symbolizer.h | 4 +
.../sanitizer_symbolizer_libcdep.cpp | 15 +
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 +
56 files changed, 3031 insertions(+), 5 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 e1abc7eb5c44e9..02ba1edb2de155 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.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_offload.cpp
index 57255546ea8885..adb515e38b848f 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_offload.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_offload.cpp
@@ -102,7 +102,7 @@ bool Offload::Discover() {
Out) == HSA_STATUS_SUCCESS;
};
auto Pool = [&](hsa_agent_t Agent) {
- hsa_amd_memory_pool_t Found{};
+ hsa_amd_memory_pool_t Fine{};
Iterate<hsa_amd_memory_pool_t>(
Api.hsa_amd_agent_iterate_memory_pools, Agent,
[&](hsa_amd_memory_pool_t Mem) {
@@ -118,10 +118,10 @@ bool Offload::Discover() {
HSA_STATUS_SUCCESS)
return HSA_STATUS_SUCCESS;
if (Flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED)
- Found = Mem;
+ Fine = Mem;
return HSA_STATUS_SUCCESS;
});
- return Found;
+ return Fine;
};
CheckHsa(Iterate<hsa_agent_t>(Api.hsa_iterate_agents, [&](hsa_agent_t Agent) {
@@ -296,6 +296,49 @@ bool Offload::Alloc(const Device& D, uptr Bytes, void** Out) {
return true;
}
+bool Offload::GetMemoryPool(hsa_agent_t Agent, hsa_amd_memory_pool_t* Pool) {
+ Lock L(&OffloadMtx);
+ if (!Ready() || !Pool)
+ return false;
+ *Pool = {};
+ hsa_status_t Status = Iterate<hsa_amd_memory_pool_t>(
+ Api.hsa_amd_agent_iterate_memory_pools, Agent,
+ [&](hsa_amd_memory_pool_t Mem) {
+ hsa_amd_segment_t Segment;
+ u32 Flags = 0;
+ bool Allowed = false;
+ if (Api.hsa_amd_memory_pool_get_info(Mem,
+ HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+ &Segment) != HSA_STATUS_SUCCESS ||
+ Segment != HSA_AMD_SEGMENT_GLOBAL ||
+ Api.hsa_amd_memory_pool_get_info(
+ Mem, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &Flags) !=
+ HSA_STATUS_SUCCESS ||
+ !(Flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) ||
+ Api.hsa_amd_memory_pool_get_info(
+ Mem, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED,
+ &Allowed) != HSA_STATUS_SUCCESS ||
+ !Allowed)
+ return HSA_STATUS_SUCCESS;
+ *Pool = Mem;
+ return HSA_STATUS_SUCCESS;
+ });
+ return Status == HSA_STATUS_SUCCESS && Pool->handle;
+}
+
+bool Offload::Allocate(hsa_amd_memory_pool_t Pool, uptr Bytes, void** Out) {
+ Lock L(&OffloadMtx);
+ if (!Ready() || !Pool.handle || !Bytes || !Out)
+ return false;
+ void* P = nullptr;
+ if (Api.hsa_amd_memory_pool_allocate(Pool, Bytes, 0, &P) !=
+ HSA_STATUS_SUCCESS ||
+ !P)
+ return false;
+ *Out = P;
+ return true;
+}
+
void Offload::Free(void* P) { Api.hsa_amd_memory_pool_free(P); }
bool Offload::Copy(void* Dst, const void* Src, uptr N) {
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_offload.h b/compiler-rt/lib/sanitizer_common/sanitizer_offload.h
index 4df2a2ca11f054..5a39e4103546c1 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_offload.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_offload.h
@@ -40,7 +40,10 @@ class Offload {
void TrackExecutable(hsa_executable_t Exec);
void UntrackExecutable(hsa_executable_t Exec);
void UntrackImages();
+ bool GetMemoryPool(hsa_agent_t Agent, hsa_amd_memory_pool_t* Pool);
+ bool Allocate(hsa_amd_memory_pool_t Pool, uptr Bytes, void** Out);
SymbolizedStack* Symbolize(uptr PC);
+ bool SymbolizeData(uptr Addr, DataInfo* Info);
private:
friend struct OffloadRpc;
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_offload_hsa.h b/compiler-rt/lib/sanitizer_common/sanitizer_offload_hsa.h
index 2300f9fa2bfabb..d2a92832539f2d 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_offload_hsa.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_offload_hsa.h
@@ -78,10 +78,12 @@ typedef enum { HSA_AMD_SEGMENT_GLOBAL = 0 } hsa_amd_segment_t;
typedef enum {
HSA_AMD_MEMORY_POOL_INFO_SEGMENT = 0,
HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS = 1,
+ HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED = 5,
} hsa_amd_memory_pool_info_t;
typedef enum {
HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED = 2,
+ HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED = 4,
} hsa_amd_memory_pool_global_flag_t;
typedef enum {
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_offload_image.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_offload_image.cpp
index eb0c5a468c5f24..09c54f7ea5c06f 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_offload_image.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_offload_image.cpp
@@ -145,4 +145,25 @@ SymbolizedStack* Offload::Symbolize(uptr PC) {
return Frames;
}
+bool Offload::SymbolizeData(uptr Addr, DataInfo* Info) {
+ if (!Addr || !Info)
+ return false;
+
+ char* Path = nullptr;
+ uptr Offset = 0;
+ if (!SnapshotImage(Addr, &Path, &Offset) || !Path)
+ return false;
+
+ bool Symbolized =
+ Symbolizer::GetOrInit()->SymbolizeModuleData(Path, Offset, Info);
+ InternalFree(Path);
+ if (!Symbolized)
+ return false;
+ if (!Info->name || !Info->name[0] || Info->name[0] == '?') {
+ Info->Clear();
+ return false;
+ }
+ return true;
+}
+
} // 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/lib/sanitizer_common/sanitizer_symbolizer.h b/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer.h
index d2aa2288f483cd..d40df017bdda63 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer.h
@@ -143,6 +143,10 @@ class Symbolizer final {
// current executing process, such as an offloading device.
SymbolizedStack* SymbolizeModuleOffset(const char* module_name,
uptr module_offset);
+ // Like SymbolizeData, but the module is not mapped in this process
+ // (offload device images).
+ bool SymbolizeModuleData(const char* module_name, uptr module_offset,
+ DataInfo* info);
bool SymbolizeData(uptr address, DataInfo *info);
bool SymbolizeFrame(uptr address, FrameInfo *info);
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer_libcdep.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer_libcdep.cpp
index b5a85c05efde7d..8ba6a14b420cb4 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer_libcdep.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_symbolizer_libcdep.cpp
@@ -114,6 +114,21 @@ SymbolizedStack* Symbolizer::SymbolizeModuleOffset(const char* module_name,
return res;
}
+bool Symbolizer::SymbolizeModuleData(const char* module_name,
+ uptr module_offset, DataInfo* info) {
+ Lock l(&mu_);
+ info->Clear();
+ info->module = internal_strdup(module_name);
+ info->module_offset = module_offset;
+ info->module_arch = kModuleArchUnknown;
+ for (auto& tool : tools_) {
+ SymbolizerScope sym_scope(this);
+ if (tool.SymbolizeData(module_offset, info))
+ return true;
+ }
+ return false;
+}
+
bool Symbolizer::SymbolizeData(uptr addr, DataInfo *info) {
Lock l(&mu_);
const char *module_name = nullptr;
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:
More information about the llvm-branch-commits
mailing list