[compiler-rt] [llvm] [compiler-rt] Add KCSAN inspired GPU ConcurrencySanitizer runtime (PR #207714)
Joseph Huber via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 14 20:47:30 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/207714
>From 94b09dad33391eb5c334ae363be612edbda1279c Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Wed, 10 Jun 2026 08:19:47 -0500
Subject: [PATCH 1/2] [compiler-rt] Add KCSAN inspired GPU ConcurrencySanitizer
runtime
Summary:
Add a standalone GPU concurrency sanitizer runtime as a new top-level
compiler-rt component 'csan' (libclang_rt.csan.a), enabled for offloading
device builds. It is a watchpoint-based data race detector modelled on the
Linux kernel's KCSAN and reports races to the host over RPC. Tests live in
the top-level 'test/csan' suite.
---
compiler-rt/cmake/caches/AMDGPU.cmake | 2 +-
compiler-rt/cmake/config-ix.cmake | 8 +-
compiler-rt/include/CMakeLists.txt | 1 +
compiler-rt/include/sanitizer/gpu_sanitizer.h | 59 ++
compiler-rt/lib/csan/.clang-format | 3 +
compiler-rt/lib/csan/CMakeLists.txt | 31 +
compiler-rt/lib/csan/csan_gpu.cpp | 732 ++++++++++++++++++
compiler-rt/test/csan/CMakeLists.txt | 29 +
compiler-rt/test/csan/array_race.c | 18 +
compiler-rt/test/csan/atomic.c | 9 +
compiler-rt/test/csan/disjoint.c | 12 +
compiler-rt/test/csan/global_race.c | 14 +
compiler-rt/test/csan/lds_disjoint.c | 12 +
compiler-rt/test/csan/lds_race.c | 15 +
compiler-rt/test/csan/lit.cfg.py | 34 +
compiler-rt/test/csan/lit.site.cfg.py.in | 11 +
compiler-rt/test/csan/memcpy_large_race.c | 25 +
compiler-rt/test/csan/memcpy_race.c | 16 +
compiler-rt/test/csan/memmove_race.c | 25 +
compiler-rt/test/csan/race.h | 28 +
compiler-rt/test/csan/single_thread.c | 9 +
compiler-rt/test/csan/write_read_race.c | 20 +
offload/plugins-nextgen/CMakeLists.txt | 2 +-
offload/plugins-nextgen/common/CMakeLists.txt | 3 +
.../common/include/PluginInterface.h | 2 +
offload/plugins-nextgen/common/include/RPC.h | 15 +
.../common/include/Sanitizer.h | 42 +
.../common/include/Utils/ELF.h | 33 +
offload/plugins-nextgen/common/src/RPC.cpp | 16 +-
.../plugins-nextgen/common/src/Sanitizer.cpp | 191 +++++
.../plugins-nextgen/common/src/Utils/ELF.cpp | 55 ++
31 files changed, 1468 insertions(+), 4 deletions(-)
create mode 100644 compiler-rt/include/sanitizer/gpu_sanitizer.h
create mode 100644 compiler-rt/lib/csan/.clang-format
create mode 100644 compiler-rt/lib/csan/CMakeLists.txt
create mode 100644 compiler-rt/lib/csan/csan_gpu.cpp
create mode 100644 compiler-rt/test/csan/CMakeLists.txt
create mode 100644 compiler-rt/test/csan/array_race.c
create mode 100644 compiler-rt/test/csan/atomic.c
create mode 100644 compiler-rt/test/csan/disjoint.c
create mode 100644 compiler-rt/test/csan/global_race.c
create mode 100644 compiler-rt/test/csan/lds_disjoint.c
create mode 100644 compiler-rt/test/csan/lds_race.c
create mode 100644 compiler-rt/test/csan/lit.cfg.py
create mode 100644 compiler-rt/test/csan/lit.site.cfg.py.in
create mode 100644 compiler-rt/test/csan/memcpy_large_race.c
create mode 100644 compiler-rt/test/csan/memcpy_race.c
create mode 100644 compiler-rt/test/csan/memmove_race.c
create mode 100644 compiler-rt/test/csan/race.h
create mode 100644 compiler-rt/test/csan/single_thread.c
create mode 100644 compiler-rt/test/csan/write_read_race.c
create mode 100644 offload/plugins-nextgen/common/include/Sanitizer.h
create mode 100644 offload/plugins-nextgen/common/src/Sanitizer.cpp
diff --git a/compiler-rt/cmake/caches/AMDGPU.cmake b/compiler-rt/cmake/caches/AMDGPU.cmake
index f3a9510c4f311..f9b4be733a5bc 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_minimal" CACHE STRING "")
+set(COMPILER_RT_SANITIZERS_TO_BUILD "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 083f1c98d0f16..43fa8aa434321 100644
--- a/compiler-rt/cmake/config-ix.cmake
+++ b/compiler-rt/cmake/config-ix.cmake
@@ -765,7 +765,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}")
@@ -880,6 +880,12 @@ else()
set(COMPILER_RT_HAS_TSAN FALSE)
endif()
+if(COMPILER_RT_GPU_BUILD)
+ set(COMPILER_RT_HAS_CSAN TRUE)
+else()
+ set(COMPILER_RT_HAS_CSAN FALSE)
+endif()
+
if (OS_NAME MATCHES "Linux|FreeBSD|Windows|NetBSD|SunOS")
set(COMPILER_RT_TSAN_HAS_STATIC_RUNTIME TRUE)
else()
diff --git a/compiler-rt/include/CMakeLists.txt b/compiler-rt/include/CMakeLists.txt
index eb998478b081b..05c78eeace345 100644
--- a/compiler-rt/include/CMakeLists.txt
+++ b/compiler-rt/include/CMakeLists.txt
@@ -5,6 +5,7 @@ if (COMPILER_RT_BUILD_SANITIZERS)
sanitizer/common_interface_defs.h
sanitizer/coverage_interface.h
sanitizer/dfsan_interface.h
+ sanitizer/gpu_sanitizer.h
sanitizer/hwasan_interface.h
sanitizer/linux_syscall_hooks.h
sanitizer/lsan_interface.h
diff --git a/compiler-rt/include/sanitizer/gpu_sanitizer.h b/compiler-rt/include/sanitizer/gpu_sanitizer.h
new file mode 100644
index 0000000000000..e38c677953c62
--- /dev/null
+++ b/compiler-rt/include/sanitizer/gpu_sanitizer.h
@@ -0,0 +1,59 @@
+//===-- sanitizer/gpu_sanitizer.h -------------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Shared ABI contract for GPU sanitizer reports shipped to the host over RPC.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef SANITIZER_GPU_SANITIZER_H
+#define SANITIZER_GPU_SANITIZER_H
+
+#include <stdint.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define SANITIZER_GPU_OPCODE(n) (('s' << 24) | (n))
+
+#define TSAN_GPU_REPORT_OPCODE SANITIZER_GPU_OPCODE(1)
+
+/// Access classification flags, mirroring the Linux kernel's KCSAN model.
+enum {
+ TSAN_GPU_ACCESS_WRITE = 1 << 0, ///< Non-atomic write.
+ TSAN_GPU_ACCESS_COMPOUND = 1 << 1, ///< Read-modify-write.
+ TSAN_GPU_ACCESS_ATOMIC = 1 << 2, ///< Atomic access.
+};
+
+/// Race report kinds.
+enum {
+ TSAN_GPU_DATA_RACE = 0, ///< Conflicting access.
+ TSAN_GPU_UNKNOWN_ORIGIN = 1, ///< Watched value changed with no finder.
+ TSAN_GPU_INTRA_WAVE = 2, ///< Conflict between lanes of the same wave.
+};
+
+/// The data associated with a single detected race.
+typedef struct __tsan_gpu_race {
+ uint64_t pc; ///< Device PC of the reporting access.
+ uint64_t peer_pc; ///< Device PC of the conflicting access.
+ uint64_t addr; ///< Accessed device address.
+ uint32_t size; ///< Access size in bytes.
+ uint32_t access_type; ///< Bitwise-or of TSAN_GPU_ACCESS_* flags.
+ uint32_t kind; ///< One of the TSAN_GPU_* race kinds.
+ uint32_t block[3]; ///< Block / workgroup id.
+ uint16_t thread[3]; ///< Thread / work-item id within the block.
+ uint8_t lane; ///< Lane id within the wave.
+ uint8_t peer_lane; ///< Conflicting lane for intra-wave races.
+ uint16_t peer_thread[3]; ///< Conflicting thread id for intra-wave races.
+} __tsan_gpu_race;
+
+#ifdef __cplusplus
+} // extern "C"
+#endif
+
+#endif // SANITIZER_GPU_SANITIZER_H
diff --git a/compiler-rt/lib/csan/.clang-format b/compiler-rt/lib/csan/.clang-format
new file mode 100644
index 0000000000000..1f2a97030379d
--- /dev/null
+++ b/compiler-rt/lib/csan/.clang-format
@@ -0,0 +1,3 @@
+BasedOnStyle: Google
+AllowShortIfStatementsOnASingleLine: false
+IndentPPDirectives: AfterHash
diff --git a/compiler-rt/lib/csan/CMakeLists.txt b/compiler-rt/lib/csan/CMakeLists.txt
new file mode 100644
index 0000000000000..646789ec317f6
--- /dev/null
+++ b/compiler-rt/lib/csan/CMakeLists.txt
@@ -0,0 +1,31 @@
+set(CSAN_SOURCES
+ csan_gpu.cpp
+ )
+
+include_directories(..)
+
+include_directories(${COMPILER_RT_SOURCE_DIR}/include)
+include(FindLibcCommonUtils)
+get_target_property(LIBC_COMMON_UTILS_INCLUDE
+ llvm-libc-common-utilities INTERFACE_INCLUDE_DIRECTORIES)
+include_directories(${LIBC_COMMON_UTILS_INCLUDE})
+
+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_object_libraries(RTCsan
+ OS ${SANITIZER_COMMON_SUPPORTED_OS}
+ ARCHS ${UBSAN_COMMON_SUPPORTED_ARCH}
+ SOURCES ${CSAN_SOURCES} CFLAGS ${CSAN_CFLAGS})
+
+add_compiler_rt_runtime(clang_rt.csan
+ STATIC
+ OS ${UBSAN_SUPPORTED_OS}
+ ARCHS ${UBSAN_SUPPORTED_ARCH}
+ OBJECT_LIBS RTCsan
+ CFLAGS ${CSAN_CFLAGS}
+ PARENT_TARGET csan)
diff --git a/compiler-rt/lib/csan/csan_gpu.cpp b/compiler-rt/lib/csan/csan_gpu.cpp
new file mode 100644
index 0000000000000..27056be71fa3a
--- /dev/null
+++ b/compiler-rt/lib/csan/csan_gpu.cpp
@@ -0,0 +1,732 @@
+//===-- csan_gpu.cpp - GPU ConcurrencySanitizer runtime -------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Watchpoint-based data race detector for GPU targets, modelled on the Linux
+// kernel's KCSAN.
+//
+//===----------------------------------------------------------------------===//
+
+#include <gpuintrin.h>
+#include <stdint.h>
+
+#include "sanitizer/gpu_sanitizer.h"
+#include "sanitizer_common/sanitizer_internal_defs.h"
+#include "shared/rpc.h"
+
+[[gnu::visibility("protected"),
+ gnu::weak]] rpc::Client client asm("__llvm_rpc_client");
+
+// Running total of races this device has reported.
+extern "C" SANITIZER_INTERFACE_ATTRIBUTE uint64_t __tsan_num_data_races = 0;
+
+// 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 uint64_t seen[64] = {};
+ const uint64_t token = (reinterpret_cast<uintptr_t>(pc) >> 4) ^
+ (static_cast<uint64_t>(kind) * 0x9E3779B97F4A7C15ull);
+ uint64_t idx = (token * 0x9E3779B97F4A7C15ull) >> 58;
+ uint64_t last = __scoped_atomic_exchange_n(
+ &seen[idx], token, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
+ return last != token;
+}
+
+// Report a data race to the RPC server so it can be symbolized and presented.
+[[gnu::cold]] static void report(unsigned kind, uintptr_t addr, uint32_t size,
+ int access_type, void* pc,
+ void* peer = nullptr, uint8_t peer_lane = 0,
+ const uint16_t (&peer_thread)[3] = {}) {
+ if (!should_report(pc, kind))
+ return;
+
+ __tsan_gpu_race rep;
+ rep.pc = reinterpret_cast<uintptr_t>(pc);
+ rep.peer_pc = reinterpret_cast<uintptr_t>(peer);
+ rep.addr = addr;
+ rep.size = size;
+ rep.access_type = static_cast<unsigned>(access_type);
+ 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_thread[0] = peer_thread ? peer_thread[0] : 0;
+ rep.peer_thread[1] = peer_thread ? peer_thread[1] : 0;
+ rep.peer_thread[2] = peer_thread ? peer_thread[2] : 0;
+ rep.kind = kind;
+
+ rpc::Client::Port Port = client.open<TSAN_GPU_REPORT_OPCODE>();
+ Port.send([&](rpc::Buffer* buf, uint32_t) {
+ __builtin_memcpy(buf->data, &rep, sizeof(rep));
+ });
+ static_assert(sizeof(__tsan_gpu_race) <= sizeof(rpc::Buffer),
+ "Report must fit in a single packet");
+
+ __scoped_atomic_fetch_add(&__tsan_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 uint64_t CLOCK_FREQ_HZ = 100000000UL;
+#else
+static constexpr uint64_t CLOCK_FREQ_HZ = 1000000000UL;
+#endif
+static constexpr uint64_t TICKS_PER_SEC = 1000000000UL;
+
+// Randomized bounds, in nanoseconds, for the watchpoint stall window. A wider
+// window catches more concurrent races at the cost of extra runtime.
+static constexpr uint64_t SAMPLE_DELAY_MIN_NS = 1000;
+static constexpr uint64_t SAMPLE_DELAY_MAX_NS = 5000;
+
+// Watchpoint table capacity, in bytes.
+static constexpr uint64_t WP_TABLE_SIZE = /*2 MiB=*/2 * 1024ul * 1024ul;
+
+// Number of slots in the watchpoint table, must be a power of two;
+static constexpr uint64_t WP_TABLE_SLOTS = WP_TABLE_SIZE / sizeof(uint64_t);
+
+// The probability that we set up a watchpoint at any given access.
+static constexpr uint32_t WP_CHANCE = 8;
+
+// Whether global accesses use the watchpoint table.
+static constexpr bool WP_ENABLE_TABLE = true;
+
+// The largest size we encode. We check a very narrow range to widen
+// watchpoints. Wide accesses are relatively uncommon on global memory.
+static constexpr uint32_t WP_MAX_SIZE = 8;
+
+static constexpr uint32_t WP_SIZE_BITS =
+ __builtin_popcountg(WP_MAX_SIZE - 1u) + 1;
+static constexpr uint32_t WP_ADDR_BITS = 64 - 2 - WP_SIZE_BITS;
+
+//===----------------------------------------------------------------------===//
+// Watchpoint encoding:
+// [63] is_write
+// [62] consumed
+// [61 : WP_ADDR_BITS] size
+// [WP_ADDR_BITS-1 : 0] address
+//===----------------------------------------------------------------------===//
+
+static constexpr uint64_t WP_INVALID = 0;
+static constexpr uint64_t WP_CONSUMED_MASK = 1ull << 62;
+static constexpr uint64_t WP_WRITE_MASK = 1ull << 63;
+static constexpr uint64_t WP_ADDR_MASK = (1ull << WP_ADDR_BITS) - 1;
+static constexpr uint64_t WP_SIZE_MASK = ((1ull << WP_SIZE_BITS) - 1)
+ << WP_ADDR_BITS;
+static_assert((WP_MAX_SIZE & (WP_MAX_SIZE - 1)) == 0,
+ "WP_MAX_SIZE must be a power of two");
+static_assert((WP_WRITE_MASK ^ WP_CONSUMED_MASK ^ WP_SIZE_MASK ^
+ WP_ADDR_MASK) == ~0ull,
+ "watchpoint fields must partition the 64-bit word");
+
+static constexpr uint64_t encode_watchpoint(uint64_t addr, uint32_t size,
+ bool is_write) {
+ return (is_write ? WP_WRITE_MASK : 0) | ((uint64_t)size << WP_ADDR_BITS) |
+ (addr & WP_ADDR_MASK);
+}
+
+static constexpr bool decode_watchpoint(uint64_t wp, uint64_t& addr,
+ uint32_t& size, bool& is_write) {
+ if (wp == WP_INVALID || (wp & WP_CONSUMED_MASK))
+ return false;
+ is_write = (wp & WP_WRITE_MASK) != 0;
+ size = (wp & WP_SIZE_MASK) >> WP_ADDR_BITS;
+ addr = wp & WP_ADDR_MASK;
+ return true;
+}
+
+static constexpr bool ranges_overlap(uint64_t a1, uint32_t s1, uint64_t a2,
+ uint32_t s2) {
+ return a1 < a2 + s2 && a2 < a1 + s1;
+}
+
+static_assert((WP_TABLE_SLOTS & (WP_TABLE_SLOTS - 1)) == 0,
+ "WP_TABLE_SLOTS must be a power of two");
+
+static constexpr uint32_t watchpoint_slot(uint64_t addr) {
+ const uint64_t word = addr >> (WP_SIZE_BITS - 1);
+ return word & (WP_TABLE_SLOTS - 1);
+}
+
+static inline uint32_t xorshift32(uint32_t& state) {
+ state ^= state << 13;
+ state ^= state >> 17;
+ state ^= state << 5;
+ return state * 0x9e3779bb;
+}
+
+// Wave-uniform Bernoulli trial that is true with probability 1 / N.
+template <uint32_t N>
+static bool bernoulli(uint32_t& rand) {
+ static_assert((N & (N - 1)) == 0,
+ "sample denominator must be a power of two");
+ [[maybe_unused]] const uint32_t r = xorshift32(rand);
+ if constexpr (N == 0)
+ return false;
+ else if constexpr (N == 1)
+ return true;
+ else
+ return (r >> (32 - __builtin_ctzg(N))) == 0;
+}
+
+static uint64_t watchpoints[WP_TABLE_SLOTS] = {};
+
+// Flatten the block-local thread index into a linear lane index.
+static uint32_t flat_thread_id() {
+ return __gpu_thread_id(__GPU_X_DIM) +
+ __gpu_num_threads(__GPU_X_DIM) *
+ (__gpu_thread_id(__GPU_Y_DIM) +
+ __gpu_num_threads(__GPU_Y_DIM) * __gpu_thread_id(__GPU_Z_DIM));
+}
+
+// PRNG seed based off of the current thread's IDs and clock cycle.
+static uint32_t entropy() {
+ return (static_cast<uint32_t>(__builtin_readcyclecounter() | 1u) ^
+ flat_thread_id() ^ (__gpu_block_id(__GPU_X_DIM) * 0x9e3779b9u));
+}
+
+namespace {
+// Per-wavefront analysis state, seeded once at kernel entry and carried for the
+// lifetime of the wave.
+struct Ctx {
+ uint32_t rand;
+};
+} // namespace
+
+// Return the calling wavefront's context from the global LDS.
+static __gpu_local Ctx& get_ctx() {
+ static constexpr uint32_t MAX_WAVEFRONTS = 1024 / 32;
+ static __gpu_local Ctx ctx[MAX_WAVEFRONTS]
+ __attribute__((loader_uninitialized));
+ return ctx[flat_thread_id() / __gpu_num_lanes()];
+}
+
+// Type trait helpers for address spaces.
+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;
+};
+
+template <typename PtrTy>
+static constexpr bool uses_table() {
+ return WP_ENABLE_TABLE && !is_ptr_local<PtrTy>::value;
+}
+
+// Returns if a wavefront should sample this access for detection. Each eligible
+// access is sampled with a wave-uniform probability of 1/SAMPLE_CHANCE.
+static bool should_watch(uint64_t lane_mask, uint32_t access_type) {
+ // If every access is atomic we cannot have a race.
+ if (!__gpu_ballot(lane_mask, !(access_type & TSAN_GPU_ACCESS_ATOMIC)))
+ return false;
+
+ bool sample = false;
+ if (__gpu_is_first_in_lane(lane_mask))
+ sample = bernoulli<WP_CHANCE>(get_ctx().rand);
+ return __gpu_read_first_lane_u32(lane_mask, sample);
+}
+
+// Scan the watchpoint table to see if there are any points set for this
+// address. Cheap lookup done for each lane in the wavefront.
+template <typename PtrTy>
+static uint64_t* find_watchpoint(uintptr_t addr, uint32_t size,
+ bool expect_write, uint64_t& encoded) {
+ // Local addresses never use the global watchpoint table.
+ if constexpr (!uses_table<PtrTy>())
+ return nullptr;
+
+ uint64_t* wp = &watchpoints[watchpoint_slot(addr)];
+ encoded = __scoped_atomic_load_n(wp, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE);
+
+ uint64_t wp_addr;
+ uint32_t wp_size;
+ bool is_write;
+ if (!decode_watchpoint(encoded, wp_addr, wp_size, is_write))
+ return nullptr;
+ if (expect_write && !is_write)
+ return nullptr;
+ if (ranges_overlap(wp_addr, wp_size, addr & WP_ADDR_MASK, size))
+ return wp;
+ return nullptr;
+}
+
+// Claim a free slot for this access by atomically flipping it from INVALID to
+// the encoded watchpoint. Returns nullptr if every candidate slot is taken.
+static uint64_t* insert_watchpoint(uint64_t addr, uint32_t size,
+ bool is_write) {
+ uint64_t* wp = &watchpoints[watchpoint_slot(addr)];
+ const uint64_t encoded = encode_watchpoint(addr, size, is_write);
+ uint64_t expected = WP_INVALID;
+ if (__scoped_atomic_compare_exchange_n(wp, &expected, encoded, false,
+ __ATOMIC_RELEASE, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE))
+ return wp;
+ return nullptr;
+}
+
+// A finder consumes a watchpoint to signal the setter that a race was observed,
+// storing its own PC into the address bits so the setter can report it as the
+// peer. Succeeds only if the slot still holds the value the finder matched.
+static bool try_consume_watchpoint(uint64_t* wp, uint64_t encoded, void* pc) {
+ uint64_t consumed =
+ WP_CONSUMED_MASK | (reinterpret_cast<uintptr_t>(pc) & WP_ADDR_MASK);
+ return __scoped_atomic_compare_exchange_n(wp, &encoded, consumed, false,
+ __ATOMIC_RELEASE, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+}
+
+// A setter tears down its own watchpoint and checks if another thread consumed
+// it and triggered a race, recovering the finder's stashed PC as 'peer'.
+static bool consume_watchpoint(uint64_t* wp, void*& peer) {
+ uint64_t old = __scoped_atomic_exchange_n(
+ wp, WP_CONSUMED_MASK, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE);
+ peer = reinterpret_cast<void*>(old & WP_ADDR_MASK);
+ return !(old & WP_CONSUMED_MASK);
+}
+
+// Release the slot back to the pool so it may be reused immediately after.
+static void remove_watchpoint(uint64_t* wp) {
+ __scoped_atomic_store_n(wp, WP_INVALID, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_DEVICE);
+}
+
+// FNV-1a digest of a byte range so wide accesses can reuse the 64-bit value
+// comparison semantics.
+template <typename BytePtr, typename WordPtr>
+static uint64_t read_range(BytePtr bytes, [[maybe_unused]] WordPtr words,
+ uint32_t size, int scope) {
+ uint64_t sum = 0xcbf29ce484222325ull;
+ uint32_t i = 0;
+
+ // Digest the unaligned prefix byte-by-byte up to a word.
+ for (; i < size && ((reinterpret_cast<uintptr_t>(bytes) + i) & 7u); ++i)
+ sum = (sum ^ __scoped_atomic_load_n(bytes + i, __ATOMIC_RELAXED, scope)) *
+ 0x100000001b3ull;
+
+ // Digest the aligned interior with wide loads.
+ for (; i + 8 <= size; i += 8)
+ sum = (sum ^ __scoped_atomic_load_n(reinterpret_cast<WordPtr>(bytes + i),
+ __ATOMIC_RELAXED, scope)) *
+ 0x100000001b3ull;
+
+ // Digest the trailing bytes that do not fill a word.
+ for (; i < size; ++i)
+ sum = (sum ^ __scoped_atomic_load_n(bytes + i, __ATOMIC_RELAXED, scope)) *
+ 0x100000001b3ull;
+ return sum;
+}
+
+// Snapshot the watched location for value-change detection. Larger sizes get
+// converted into a single checksum.
+static uint64_t read_instrumented_memory(const volatile __gpu_global void* ptr,
+ uint32_t size) {
+ const uintptr_t addr =
+ reinterpret_cast<uintptr_t>(const_cast<const __gpu_global void*>(ptr));
+ if ((addr & (size - 1)) == 0) {
+ switch (size) {
+ case 1:
+ return __scoped_atomic_load_n((const volatile __gpu_global uint8_t*)ptr,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+ case 2:
+ return __scoped_atomic_load_n(
+ (const volatile __gpu_global uint16_t*)ptr, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_SYSTEM);
+ case 4:
+ return __scoped_atomic_load_n(
+ (const volatile __gpu_global uint32_t*)ptr, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_SYSTEM);
+ case 8:
+ return __scoped_atomic_load_n(
+ (const volatile __gpu_global uint64_t*)ptr, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_SYSTEM);
+ }
+ }
+ return read_range((const volatile __gpu_global uint8_t*)ptr,
+ (const volatile __gpu_global uint64_t*)ptr, size,
+ __MEMORY_SCOPE_SYSTEM);
+}
+
+static uint64_t read_instrumented_memory(const volatile __gpu_local void* ptr,
+ uint32_t size) {
+ // LDS is only coherent within the workgroup, so the scope narrows to match.
+ const uintptr_t addr =
+ reinterpret_cast<uintptr_t>(const_cast<const __gpu_local void*>(ptr));
+ if ((addr & (size - 1)) == 0) {
+ switch (size) {
+ case 1:
+ return __scoped_atomic_load_n((const volatile __gpu_local uint8_t*)ptr,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP);
+ case 2:
+ return __scoped_atomic_load_n((const volatile __gpu_local uint16_t*)ptr,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP);
+ case 4:
+ return __scoped_atomic_load_n((const volatile __gpu_local uint32_t*)ptr,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP);
+ case 8:
+ return __scoped_atomic_load_n((const volatile __gpu_local uint64_t*)ptr,
+ __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP);
+ }
+ }
+ return read_range((const volatile __gpu_local uint8_t*)ptr,
+ (const volatile __gpu_local uint64_t*)ptr, size,
+ __MEMORY_SCOPE_WRKGRP);
+}
+
+// A conflicting watchpoint already exists. Try to consume it so the thread that
+// created the watchpoint will see the conflict.
+template <typename PtrTy>
+[[gnu::cold]] static void found_watchpoint(const PtrTy addr, uint32_t size,
+ int access_type, void* pc,
+ uint64_t* wp, uint64_t encoded) {
+ try_consume_watchpoint(wp, encoded, pc);
+}
+
+// Trivial check to catch races between lanes inside of the same wave.
+static bool intra_wave_race(uint64_t lane_mask, uintptr_t addr, int access_type,
+ uint8_t& peer_lane, uint16_t peer_thread[3]) {
+ const bool is_write = (access_type & TSAN_GPU_ACCESS_WRITE) != 0;
+ const bool is_atomic = (access_type & TSAN_GPU_ACCESS_ATOMIC) != 0;
+
+ const uint64_t same_addr = __gpu_match_any_u64(lane_mask, addr);
+ const uint64_t writers = __gpu_ballot(lane_mask, is_write);
+ const uint64_t nonatomic = __gpu_ballot(lane_mask, !is_atomic);
+
+ // Every lane reads the peer's coordinates so the source lane is always live.
+ const uint32_t peer = 63u - __builtin_clzg(same_addr);
+ const uint32_t width = __gpu_num_lanes();
+ peer_lane = static_cast<uint8_t>(peer);
+ peer_thread[0] = __gpu_shuffle_idx_u32(lane_mask, peer,
+ __gpu_thread_id(__GPU_X_DIM), width);
+ peer_thread[1] = __gpu_shuffle_idx_u32(lane_mask, peer,
+ __gpu_thread_id(__GPU_Y_DIM), width);
+ peer_thread[2] = __gpu_shuffle_idx_u32(lane_mask, peer,
+ __gpu_thread_id(__GPU_Z_DIM), width);
+
+ // The first lane of each racing group reports its own conflict.
+ const bool is_race = __builtin_popcountg(same_addr) >= 2 &&
+ (same_addr & writers) && (same_addr & nonatomic);
+ return is_race && __gpu_is_first_in_lane(same_addr);
+}
+
+static void delay_ns(uint64_t nsecs) {
+ const uint64_t tick_rate = TICKS_PER_SEC / CLOCK_FREQ_HZ;
+ const uint64_t start = __builtin_readsteadycounter();
+ const uint64_t 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
+}
+
+// Draw a wave-uniform random delay in the range [lo, hi) nanoseconds.
+static uint64_t random_delay(uint64_t lane_mask, uint64_t lo, uint64_t hi) {
+ uint64_t nsecs = lo;
+ if (__gpu_is_first_in_lane(lane_mask))
+ nsecs = lo + xorshift32(get_ctx().rand) % (hi - lo);
+ return __gpu_read_first_lane_u64(lane_mask, nsecs);
+}
+
+static void sample_delay(uint64_t lane_mask) {
+ delay_ns(random_delay(lane_mask, SAMPLE_DELAY_MIN_NS, SAMPLE_DELAY_MAX_NS));
+}
+
+// The slow path, sets a watchpoint in the table and waits to see if any other
+// thread tripped it while finding a watchpoint.
+template <typename PtrTy>
+static void watch(uint64_t lane_mask, const PtrTy addr, uint32_t size,
+ int access_type, void* pc) {
+ const bool is_write = (access_type & TSAN_GPU_ACCESS_WRITE) != 0;
+ const uintptr_t iaddr = reinterpret_cast<uintptr_t>(addr);
+
+ uint8_t peer_lane;
+ uint16_t peer_thread[3];
+ if (intra_wave_race(lane_mask, iaddr, access_type, peer_lane, peer_thread))
+ report(TSAN_GPU_INTRA_WAVE, iaddr, size, access_type, pc, nullptr,
+ peer_lane, peer_thread);
+
+ // Oversized accesses watch their base; atomics and LDS cannot be armed, and
+ // nothing is armed when the table is disabled (value-only mode).
+ const uint32_t wp_size = size < WP_MAX_SIZE ? size : WP_MAX_SIZE;
+ const bool armable = uses_table<PtrTy>() &&
+ !(access_type & TSAN_GPU_ACCESS_ATOMIC) && iaddr != 0;
+ uint64_t* wp =
+ armable ? insert_watchpoint(iaddr, wp_size, is_write) : nullptr;
+
+ const uint64_t old = read_instrumented_memory(addr, size);
+ sample_delay(lane_mask);
+ const uint64_t now = read_instrumented_memory(addr, size);
+
+ void* peer = nullptr;
+ if (wp && !consume_watchpoint(wp, peer))
+ // A finder consumed the watchpoint, report a race with a known peer.
+ report(TSAN_GPU_DATA_RACE, iaddr, size, access_type, pc, peer);
+ else if (old != now)
+ // The value moved under us with no finder a race of unknown origin.
+ report(TSAN_GPU_UNKNOWN_ORIGIN, iaddr, size, access_type, pc);
+
+ if (wp)
+ remove_watchpoint(wp);
+}
+
+template <typename PtrTy>
+static void check_access_impl(uint64_t lane_mask, const PtrTy addr,
+ uint32_t size, int access_type, void* pc) {
+ // Finder side, every access probes for a conflicting watchpoint.
+ if constexpr (uses_table<PtrTy>()) {
+ const bool is_write = (access_type & TSAN_GPU_ACCESS_WRITE) != 0;
+ uint64_t encoded;
+ uint64_t* wp =
+ find_watchpoint<PtrTy>((uint64_t)addr, size, !is_write, encoded);
+ if (wp)
+ found_watchpoint(addr, size, access_type, pc, wp, encoded);
+ }
+
+ // Single sampled path, decided per wave so the stall is reached in lockstep.
+ if (should_watch(lane_mask, access_type))
+ watch(lane_mask, addr, size, access_type, pc);
+}
+
+static void check_access(const volatile void* addr, uintptr_t size,
+ int access_type, void* pc) {
+ // Private addresses or empty accesses by definition cannot race.
+ if (__gpu_is_ptr_private(const_cast<void*>(addr)) || !size)
+ return;
+
+ 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);
+}
+
+// Hooks to track weak memory ordering with additional scoping information.
+// Currently unused with questionable practicality for the GPU case.
+[[maybe_unused]] static void barrier_mb(int scope) { /* TODO */ }
+[[maybe_unused]] static void barrier_wmb(int scope) { /* TODO */ }
+[[maybe_unused]] static void barrier_rmb(int scope) { /* TODO */ }
+[[maybe_unused]] static void barrier_release(int scope) { /* TODO */ }
+
+// Translate an atomic builtin memory order into the barrier it implies.
+static void atomic_memorder(int order, int scope) {
+ if (order == __ATOMIC_RELEASE || order == __ATOMIC_ACQ_REL ||
+ order == __ATOMIC_SEQ_CST)
+ barrier_release(scope);
+}
+
+// Translate a standalone fence memory order into the implied barrier.
+static void fence_memorder(int order, int scope) {
+ switch (order) {
+ case __ATOMIC_SEQ_CST:
+ case __ATOMIC_ACQ_REL:
+ barrier_mb(scope);
+ break;
+ case __ATOMIC_RELEASE:
+ barrier_release(scope);
+ break;
+ case __ATOMIC_ACQUIRE:
+ case __ATOMIC_CONSUME:
+ barrier_rmb(scope);
+ break;
+ default:
+ break;
+ }
+}
+
+// Runs at the start of every instrumented kernel to seed the per-wave context
+// once, so all later sampling and delays draw from a single entropy source.
+static void init_watchpoints() {
+ uint64_t lane_mask = __gpu_lane_mask();
+ if (__gpu_is_first_in_lane(lane_mask))
+ get_ctx().rand = entropy();
+ __gpu_sync_threads();
+}
+
+//===----------------------------------------------------------------------===//
+// Public API (ABI emitted by the ThreadSanitizer instrumentation pass)
+//===----------------------------------------------------------------------===//
+
+#define INTERFACE extern "C" SANITIZER_INTERFACE_ATTRIBUTE
+
+INTERFACE void __tsan_kernel_entry() { init_watchpoints(); }
+
+INTERFACE void __tsan_func_entry(void* pc) {}
+INTERFACE void __tsan_func_exit() {}
+
+INTERFACE void __tsan_ignore_thread_begin() {}
+INTERFACE void __tsan_ignore_thread_end() {}
+
+#define TSAN_READ(N) \
+ INTERFACE void __tsan_read##N(void* addr) { \
+ check_access(addr, N, 0, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_unaligned_read##N(void* addr) { \
+ check_access(addr, N, 0, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_volatile_read##N(void* addr) { \
+ check_access(addr, N, 0, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_unaligned_volatile_read##N(void* addr) { \
+ check_access(addr, N, 0, __builtin_return_address(0)); \
+ }
+
+#define TSAN_WRITE(N) \
+ INTERFACE void __tsan_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_unaligned_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_volatile_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_unaligned_volatile_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0)); \
+ }
+
+#define TSAN_READ_WRITE(N) \
+ INTERFACE void __tsan_read_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE | TSAN_GPU_ACCESS_COMPOUND, \
+ __builtin_return_address(0)); \
+ } \
+ INTERFACE void __tsan_unaligned_read_write##N(void* addr) { \
+ check_access(addr, N, TSAN_GPU_ACCESS_WRITE | TSAN_GPU_ACCESS_COMPOUND, \
+ __builtin_return_address(0)); \
+ }
+
+#define TSAN_ACCESS(N) \
+ TSAN_READ(N) \
+ TSAN_WRITE(N) \
+ TSAN_READ_WRITE(N)
+
+TSAN_ACCESS(1)
+TSAN_ACCESS(2)
+TSAN_ACCESS(4)
+TSAN_ACCESS(8)
+TSAN_ACCESS(16)
+
+INTERFACE void __tsan_read_range(void* addr, uintptr_t size) {
+ check_access(addr, size, 0, __builtin_return_address(0));
+}
+
+INTERFACE void __tsan_write_range(void* addr, uintptr_t size) {
+ check_access(addr, size, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0));
+}
+
+#define TSAN_ATOMIC_LOAD(N) \
+ INTERFACE uint##N##_t __tsan_atomic##N##_load(const volatile uint##N##_t* a, \
+ int order, int scope) { \
+ atomic_memorder(order, scope); \
+ check_access(a, N / 8, TSAN_GPU_ACCESS_ATOMIC, \
+ __builtin_return_address(0)); \
+ return __scoped_atomic_load_n(a, order, scope); \
+ }
+
+#define TSAN_ATOMIC_STORE(N) \
+ INTERFACE void __tsan_atomic##N##_store( \
+ volatile uint##N##_t* a, uint##N##_t v, int order, int scope) { \
+ atomic_memorder(order, scope); \
+ check_access(a, N / 8, TSAN_GPU_ACCESS_ATOMIC | TSAN_GPU_ACCESS_WRITE, \
+ __builtin_return_address(0)); \
+ __scoped_atomic_store_n(a, v, order, scope); \
+ }
+
+#define TSAN_ATOMIC_EXCHANGE(N) \
+ INTERFACE uint##N##_t __tsan_atomic##N##_exchange( \
+ volatile uint##N##_t* a, uint##N##_t v, int order, int scope) { \
+ atomic_memorder(order, scope); \
+ check_access(a, N / 8, TSAN_GPU_ACCESS_ATOMIC | TSAN_GPU_ACCESS_WRITE, \
+ __builtin_return_address(0)); \
+ return __scoped_atomic_exchange_n(a, v, order, scope); \
+ }
+
+#define TSAN_ATOMIC_FETCH_OP(N, op) \
+ INTERFACE uint##N##_t __tsan_atomic##N##_fetch_##op( \
+ volatile uint##N##_t* a, uint##N##_t v, int order, int scope) { \
+ atomic_memorder(order, scope); \
+ check_access(a, N / 8, TSAN_GPU_ACCESS_ATOMIC | TSAN_GPU_ACCESS_WRITE, \
+ __builtin_return_address(0)); \
+ return __scoped_atomic_fetch_##op(a, v, order, scope); \
+ }
+
+#define TSAN_ATOMIC_CAS(N) \
+ INTERFACE uint##N##_t __tsan_atomic##N##_compare_exchange_val( \
+ volatile uint##N##_t* a, uint##N##_t c, uint##N##_t v, int order_succ, \
+ int order_fail, int scope) { \
+ atomic_memorder(order_succ, scope); \
+ check_access(a, N / 8, TSAN_GPU_ACCESS_ATOMIC | TSAN_GPU_ACCESS_WRITE, \
+ __builtin_return_address(0)); \
+ __scoped_atomic_compare_exchange_n(a, &c, v, false, order_succ, \
+ order_fail, scope); \
+ return c; \
+ }
+
+#define TSAN_ATOMIC_OPS(N) \
+ TSAN_ATOMIC_LOAD(N) \
+ TSAN_ATOMIC_STORE(N) \
+ TSAN_ATOMIC_EXCHANGE(N) \
+ TSAN_ATOMIC_FETCH_OP(N, add) \
+ TSAN_ATOMIC_FETCH_OP(N, sub) \
+ TSAN_ATOMIC_FETCH_OP(N, and) \
+ TSAN_ATOMIC_FETCH_OP(N, or) \
+ TSAN_ATOMIC_FETCH_OP(N, xor) \
+ TSAN_ATOMIC_FETCH_OP(N, nand) \
+ TSAN_ATOMIC_CAS(N)
+
+TSAN_ATOMIC_OPS(8)
+TSAN_ATOMIC_OPS(16)
+TSAN_ATOMIC_OPS(32)
+TSAN_ATOMIC_OPS(64)
+
+INTERFACE void __tsan_atomic_thread_fence(int order, int scope) {
+ fence_memorder(order, scope);
+ __scoped_atomic_thread_fence(order, scope);
+}
+
+INTERFACE void __tsan_atomic_signal_fence(int order, int scope) {
+ fence_memorder(order, scope);
+ __atomic_signal_fence(order);
+}
+
+INTERFACE void __tsan_vptr_update(void* vptr_p, void* new_val) {}
+INTERFACE void __tsan_vptr_read(void* vptr_p) {}
+
+INTERFACE void* __tsan_memmove(void* dst, const void* src, uintptr_t sz) {
+ void* pc = __builtin_return_address(0);
+ check_access(dst, sz, TSAN_GPU_ACCESS_WRITE, pc);
+ check_access(src, sz, 0, pc);
+ return __builtin_memmove(dst, src, sz);
+}
+
+INTERFACE void* __tsan_memcpy(void* dst, const void* src, uintptr_t sz) {
+ void* pc = __builtin_return_address(0);
+ check_access(dst, sz, TSAN_GPU_ACCESS_WRITE, pc);
+ check_access(src, sz, 0, pc);
+ return __builtin_memcpy(dst, src, sz);
+}
+
+INTERFACE void* __tsan_memset(void* dst, int c, uintptr_t sz) {
+ check_access(dst, sz, TSAN_GPU_ACCESS_WRITE, __builtin_return_address(0));
+ return __builtin_memset(dst, c, sz);
+}
+
+// FIXME: Required to resolve the workgroup size without a set ABI version.
+#ifdef __AMDGPU__
+extern "C" const inline uint32_t __oclc_ABI_version = 0;
+[[gnu::alias("__oclc_ABI_version")]] const uint32_t __oclc_ABI_version__;
+#endif
diff --git a/compiler-rt/test/csan/CMakeLists.txt b/compiler-rt/test/csan/CMakeLists.txt
new file mode 100644
index 0000000000000..16ffa801cd494
--- /dev/null
+++ b/compiler-rt/test/csan/CMakeLists.txt
@@ -0,0 +1,29 @@
+set(CSAN_LIT_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
+
+if(COMPILER_RT_EMULATOR)
+ set(CSAN_TEST_DEPS ${SANITIZER_COMMON_LIT_TEST_DEPS} csan)
+ set(CSAN_TEST_TARGET_ARCH ${COMPILER_RT_DEFAULT_TARGET_ARCH})
+ set(CSAN_TEST_CONFIG_SUFFIX "-${COMPILER_RT_DEFAULT_TARGET_ARCH}")
+ get_test_cc_for_arch(${COMPILER_RT_DEFAULT_TARGET_ARCH}
+ CSAN_TEST_TARGET_CC CSAN_TEST_TARGET_CFLAGS)
+
+ if(COMPILER_RT_TARGET_AMDGPU)
+ list(APPEND CSAN_TEST_TARGET_CFLAGS
+ -Wno-multi-gpu -flto -mcpu=native -nogpulib -stdlib -startfiles)
+ else()
+ list(APPEND CSAN_TEST_TARGET_CFLAGS
+ -Wno-multi-gpu -flto -march=native -nogpulib -stdlib -startfiles)
+ endif()
+ list(APPEND CSAN_TEST_TARGET_CFLAGS
+ --target=${COMPILER_RT_DEFAULT_TARGET_TRIPLE})
+ string(REPLACE ";" " " CSAN_TEST_TARGET_CFLAGS "${CSAN_TEST_TARGET_CFLAGS}")
+
+ configure_lit_site_cfg(
+ ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
+ ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
+ MAIN_CONFIG ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py)
+ add_lit_testsuite(check-csan "Running GPU ConcurrencySanitizer tests"
+ ${CMAKE_CURRENT_BINARY_DIR}
+ DEPENDS ${CSAN_TEST_DEPS})
+ set_target_properties(check-csan PROPERTIES FOLDER "Compiler-RT Tests")
+endif()
diff --git a/compiler-rt/test/csan/array_race.c b/compiler-rt/test/csan/array_race.c
new file mode 100644
index 0000000000000..1e03d2f94527b
--- /dev/null
+++ b/compiler-rt/test/csan/array_race.c
@@ -0,0 +1,18 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include <gpuintrin.h>
+
+#include "race.h"
+
+volatile int data[64];
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: Write of size 4 at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*array_race\.c:[0-9]+:[0-9]+}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'data'
+int main(void) {
+ unsigned id = __gpu_num_threads(0) * __gpu_block_id(0) + __gpu_thread_id(0);
+ unsigned slot = id % 64;
+ RACE_UNTIL_FOUND(i) { data[slot]++; }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/atomic.c b/compiler-rt/test/csan/atomic.c
new file mode 100644
index 0000000000000..cc3cd72cd8413
--- /dev/null
+++ b/compiler-rt/test/csan/atomic.c
@@ -0,0 +1,9 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | count 0
+
+int global;
+
+int main(void) {
+ for (int i = 0; i < 1024; ++i)
+ __atomic_fetch_add(&global, 1, __ATOMIC_RELAXED);
+ return 0;
+}
diff --git a/compiler-rt/test/csan/disjoint.c b/compiler-rt/test/csan/disjoint.c
new file mode 100644
index 0000000000000..42f1da71a4980
--- /dev/null
+++ b/compiler-rt/test/csan/disjoint.c
@@ -0,0 +1,12 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | count 0
+
+#include <gpuintrin.h>
+
+static int data[64 * 64];
+
+int main(void) {
+ unsigned id = __gpu_num_threads(0) * __gpu_block_id(0) + __gpu_thread_id(0);
+ for (int i = 0; i < 1024; ++i)
+ data[id % (64 * 64)] += i;
+ return 0;
+}
diff --git a/compiler-rt/test/csan/global_race.c b/compiler-rt/test/csan/global_race.c
new file mode 100644
index 0000000000000..7689075c686fe
--- /dev/null
+++ b/compiler-rt/test/csan/global_race.c
@@ -0,0 +1,14 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+volatile int global;
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: Write of size 4 at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*global_race\.c:[0-9]+:[0-9]+}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'global'
+int main(void) {
+ RACE_UNTIL_FOUND(i) { global++; }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/lds_disjoint.c b/compiler-rt/test/csan/lds_disjoint.c
new file mode 100644
index 0000000000000..95d4c9c574e97
--- /dev/null
+++ b/compiler-rt/test/csan/lds_disjoint.c
@@ -0,0 +1,12 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | count 0
+
+#include <gpuintrin.h>
+
+[[clang::loader_uninitialized]] static __gpu_local int shared[64];
+
+int main(void) {
+ unsigned slot = __gpu_thread_id(__GPU_X_DIM) % 64;
+ for (int i = 0; i < 1024; ++i)
+ shared[slot] += i;
+ return 0;
+}
diff --git a/compiler-rt/test/csan/lds_race.c b/compiler-rt/test/csan/lds_race.c
new file mode 100644
index 0000000000000..dbf82d0ad505e
--- /dev/null
+++ b/compiler-rt/test/csan/lds_race.c
@@ -0,0 +1,15 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include <gpuintrin.h>
+
+#include "race.h"
+
+[[clang::loader_uninitialized]] static volatile __gpu_local int shared[64];
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: {{Write|Read-modify-write|Read}} of size 4 at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*lds_race\.c:[0-9]+:[0-9]+}}
+int main(void) {
+ RACE_UNTIL_FOUND(i) { shared[0] += i; }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/lit.cfg.py b/compiler-rt/test/csan/lit.cfg.py
new file mode 100644
index 0000000000000..8993c2a6502d9
--- /dev/null
+++ b/compiler-rt/test/csan/lit.cfg.py
@@ -0,0 +1,34 @@
+# -*- 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! You may need to run "
+ "tests from your build directory or add this attribute "
+ "to lit.site.cfg.py " % attr_name
+ )
+ return attr_value
+
+
+config.name = "ConcurrencySanitizer-GPU" + config.name_suffix
+config.test_source_root = os.path.dirname(__file__)
+config.suffixes = [".c", ".cpp"]
+
+if not config.emulator:
+ config.unsupported = True
+
+
+def build_invocation(compile_flags):
+ return " " + " ".join([config.clang] + compile_flags) + " "
+
+
+target_cflags = [get_required_attr(config, "target_cflags")]
+clang_csan_cflags = ["-fsanitize=concurrency", "-gline-tables-only"] + target_cflags
+clang_csan_cxxflags = config.cxx_mode_flags + clang_csan_cflags
+
+config.substitutions.append(("%clang_csan ", build_invocation(clang_csan_cflags)))
+config.substitutions.append(("%clangxx_csan ", build_invocation(clang_csan_cxxflags)))
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 0000000000000..37718a367cb8d
--- /dev/null
+++ b/compiler-rt/test/csan/lit.site.cfg.py.in
@@ -0,0 +1,11 @@
+ at LIT_SITE_CFG_IN_HEADER@
+
+config.name_suffix = "@CSAN_TEST_CONFIG_SUFFIX@"
+config.target_cflags = "@CSAN_TEST_TARGET_CFLAGS@"
+config.target_arch = "@CSAN_TEST_TARGET_ARCH@"
+
+# Load common config for all compiler-rt lit tests.
+lit_config.load_config(config, "@COMPILER_RT_BINARY_DIR@/test/lit.common.configured")
+
+# Load tool-specific config that would do the real work.
+lit_config.load_config(config, "@CSAN_LIT_SOURCE_DIR@/lit.cfg.py")
diff --git a/compiler-rt/test/csan/memcpy_large_race.c b/compiler-rt/test/csan/memcpy_large_race.c
new file mode 100644
index 0000000000000..4877b434ba3d3
--- /dev/null
+++ b/compiler-rt/test/csan/memcpy_large_race.c
@@ -0,0 +1,25 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include <gpuintrin.h>
+
+#include "race.h"
+
+#define N 512
+
+int dst[N];
+int src[N];
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: {{Write|Read}} of size {{[0-9]+}} at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*memcpy_large_race\.c:[0-9]+:[0-9]+}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'dst'
+int main(void) {
+ unsigned id = __gpu_num_threads(0) * __gpu_block_id(0) + __gpu_thread_id(0);
+ RACE_UNTIL_FOUND(i) {
+ if (id == 0)
+ __builtin_memcpy((void *)dst, (const void *)src, sizeof(dst));
+ else
+ dst[N - 1] = id;
+ }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/memcpy_race.c b/compiler-rt/test/csan/memcpy_race.c
new file mode 100644
index 0000000000000..dd48cf3116a45
--- /dev/null
+++ b/compiler-rt/test/csan/memcpy_race.c
@@ -0,0 +1,16 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+volatile unsigned len = 32;
+int dst[64];
+int src[64];
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: Write of size {{[0-9]+}} at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*memcpy_race\.c:[0-9]+:[0-9]+}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'dst'
+int main(void) {
+ RACE_UNTIL_FOUND(i) { __builtin_memcpy((void *)dst, (const void *)src, len); }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/memmove_race.c b/compiler-rt/test/csan/memmove_race.c
new file mode 100644
index 0000000000000..aa41bf1d07740
--- /dev/null
+++ b/compiler-rt/test/csan/memmove_race.c
@@ -0,0 +1,25 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include <gpuintrin.h>
+
+#include "race.h"
+
+#define N 512
+
+int buf[N];
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: {{Write|Read}} of size {{[0-9]+}} at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*memmove_race\.c:[0-9]+:[0-9]+}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'buf'
+int main(void) {
+ unsigned id = __gpu_num_threads(0) * __gpu_block_id(0) + __gpu_thread_id(0);
+ RACE_UNTIL_FOUND(i) {
+ if (id == 0)
+ __builtin_memmove((void *)buf, (const void *)(buf + 1),
+ (N - 1) * sizeof(int));
+ else
+ buf[N - 1] = id;
+ }
+ return 0;
+}
diff --git a/compiler-rt/test/csan/race.h b/compiler-rt/test/csan/race.h
new file mode 100644
index 0000000000000..df258d70d516b
--- /dev/null
+++ b/compiler-rt/test/csan/race.h
@@ -0,0 +1,28 @@
+//===-- race.h - Shared helpers for GPU ConcurrencySanitizer tests --------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TSAN_GPU_TEST_RACE_H
+#define TSAN_GPU_TEST_RACE_H
+
+// Running total of races the device runtime has reported so far.
+extern unsigned long long __tsan_num_data_races;
+
+// True once the detector has fired at least once.
+static inline int race_found(void) {
+ return __atomic_load_n(&__tsan_num_data_races, __ATOMIC_RELAXED) != 0;
+}
+
+// Timeout value so the tests do not run forever.
+#define RACE_MAX_ITERS (1 << 20)
+
+// Sanitization reports are fundamentally probabalistic. We need to sample the
+// racy region repeatedly until it fires.
+#define RACE_UNTIL_FOUND(i) \
+ for (int i = 0; i < RACE_MAX_ITERS && !race_found(); ++i)
+
+#endif // TSAN_GPU_TEST_RACE_H
diff --git a/compiler-rt/test/csan/single_thread.c b/compiler-rt/test/csan/single_thread.c
new file mode 100644
index 0000000000000..5f567fc028202
--- /dev/null
+++ b/compiler-rt/test/csan/single_thread.c
@@ -0,0 +1,9 @@
+// RUN: %clang_csan %s -o %t && %run --threads 1 --blocks 1 %t 2>&1 | count 0
+
+int global;
+
+int main(void) {
+ for (int i = 0; i < 1024; ++i)
+ global++;
+ return 0;
+}
diff --git a/compiler-rt/test/csan/write_read_race.c b/compiler-rt/test/csan/write_read_race.c
new file mode 100644
index 0000000000000..6a91d0ca6a841
--- /dev/null
+++ b/compiler-rt/test/csan/write_read_race.c
@@ -0,0 +1,20 @@
+// RUN: %clang_csan %s -o %t && %run --threads 64 --blocks 64 %t 2>&1 | FileCheck %s
+
+#include "race.h"
+
+volatile int data;
+
+// CHECK: WARNING: ConcurrencySanitizer: data race
+// CHECK: {{Write|Read}} of size 4 at 0x{{[0-9a-f]+}}
+// CHECK: #0 {{.*write_read_race\.c:[0-9]+:[0-9]+|.*\?\?}}
+// CHECK: Address 0x{{[0-9a-f]+}} is global variable 'data'
+int main(void) {
+ int sink = 0;
+ RACE_UNTIL_FOUND(i) {
+ data = i;
+ sink += data;
+ }
+ if (sink == -1)
+ __builtin_trap();
+ return 0;
+}
diff --git a/offload/plugins-nextgen/CMakeLists.txt b/offload/plugins-nextgen/CMakeLists.txt
index 78dfe7455d475..b96c96874ba8e 100644
--- a/offload/plugins-nextgen/CMakeLists.txt
+++ b/offload/plugins-nextgen/CMakeLists.txt
@@ -12,7 +12,7 @@ function(add_target_library target_name lib_name)
AggressiveInstCombine Analysis BinaryFormat BitReader BitWriter CodeGen
Core Extensions FrontendOffloading InstCombine Instrumentation IPO IRReader
Linker MC Object Passes ProfileData Remarks ScalarOpts Support Target
- TargetParser TransformUtils Vectorize)
+ TargetParser TransformUtils Vectorize DebugInfoDWARF Object)
endif()
llvm_update_compile_flags(${target_name})
target_include_directories(${target_name} PUBLIC ${common_dir}/include
diff --git a/offload/plugins-nextgen/common/CMakeLists.txt b/offload/plugins-nextgen/common/CMakeLists.txt
index 4063ce5c4e18c..f1b73cf421c9e 100644
--- a/offload/plugins-nextgen/common/CMakeLists.txt
+++ b/offload/plugins-nextgen/common/CMakeLists.txt
@@ -15,6 +15,7 @@ add_library(PluginCommon OBJECT
src/JIT.cpp
src/RecordReplay.cpp
src/RPC.cpp
+ src/Sanitizer.cpp
src/OffloadError.cpp
src/Utils/ELF.cpp
)
@@ -52,6 +53,8 @@ target_include_directories(PluginCommon PUBLIC
${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
${LIBOMPTARGET_BINARY_INCLUDE_DIR}
${LIBOMPTARGET_INCLUDE_DIR}
+ # For the shared GPU sanitizer report ABI header.
+ ${OFFLOAD_SOURCE_DIR}/../compiler-rt/include
)
set_target_properties(PluginCommon PROPERTIES POSITION_INDEPENDENT_CODE ON)
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 2a23d319fff49..ba9995a15bc95 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -16,6 +16,7 @@
#include <deque>
#include <list>
#include <map>
+#include <optional>
#include <shared_mutex>
#include <variant>
#include <vector>
@@ -41,6 +42,7 @@
#include "omp-tools.h"
#endif
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseMapInfo.h"
#include "llvm/ADT/DenseSet.h"
#include "llvm/ADT/Hashing.h"
diff --git a/offload/plugins-nextgen/common/include/RPC.h b/offload/plugins-nextgen/common/include/RPC.h
index 3175d37770577..fe206f31c19a8 100644
--- a/offload/plugins-nextgen/common/include/RPC.h
+++ b/offload/plugins-nextgen/common/include/RPC.h
@@ -24,6 +24,7 @@
#include <condition_variable>
#include <cstdint>
#include <functional>
+#include <memory>
#include <mutex>
#include <thread>
@@ -35,6 +36,10 @@ class GenericGlobalHandlerTy;
class DeviceImageTy;
} // namespace plugin
+/// Deduplication tables for GPU sanitizer diagnostics; defined in 'Sanitizer.h'
+/// and held by pointer so this header stays free of the sanitizer ABI include.
+struct SanitizerTables;
+
/// A generic class implementing the interface between the RPC server provided
/// by the 'libc' project and 'libomptarget'. If the RPC server is not available
/// these routines will perform no action.
@@ -45,6 +50,10 @@ struct RPCServerTy {
/// Initializes the handles to the number of devices we may need to service.
RPCServerTy(plugin::GenericPluginTy &Plugin);
+ /// Defined out of line so the owned 'SanitizerTables' need only be complete
+ /// in the implementation file.
+ ~RPCServerTy();
+
/// Deinitialize the associated memory and resources.
llvm::Error shutDown(plugin::GenericPluginTy &Plugin);
@@ -79,7 +88,13 @@ struct RPCServerTy {
void setSleepFunction(std::function<void()> Sleep,
std::function<void()> Wake);
+ /// Sanitizer deduplication state with the same lifetime as the RPC interface.
+ SanitizerTables &getSanitizerTables() { return *Sanitizers; }
+
private:
+ /// Deduplication tables for any sanitizer diagnostics this server services.
+ std::unique_ptr<SanitizerTables> Sanitizers;
+
/// Array from this device's identifier to its attached devices.
std::unique_ptr<void *[]> Buffers;
diff --git a/offload/plugins-nextgen/common/include/Sanitizer.h b/offload/plugins-nextgen/common/include/Sanitizer.h
new file mode 100644
index 0000000000000..fd6020afec567
--- /dev/null
+++ b/offload/plugins-nextgen/common/include/Sanitizer.h
@@ -0,0 +1,42 @@
+//===-- Sanitizer.h - Host-side GPU sanitizer reporting ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OFFLOAD_PLUGINS_NEXTGEN_COMMON_SANITIZER_H
+#define OFFLOAD_PLUGINS_NEXTGEN_COMMON_SANITIZER_H
+
+#include "sanitizer/gpu_sanitizer.h"
+
+#include "llvm/ADT/DenseSet.h"
+
+#include <cstdint>
+#include <mutex>
+#include <utility>
+
+namespace llvm::omp::target {
+namespace plugin {
+struct GenericDeviceTy;
+class DeviceImageTy;
+} // namespace plugin
+
+/// Deduplication tables for GPU sanitizer diagnostics.
+struct SanitizerTables {
+ /// Returns true the first time this conflicting PC pair and kind are seen.
+ bool isNewRace(uint64_t PC, uint64_t PeerPC, unsigned Kind);
+
+private:
+ std::mutex Mtx;
+ DenseSet<std::pair<uint64_t, uint64_t>> Races;
+};
+
+/// Report a concurrency sanitizer hit on the given device with deduplication.
+void reportGPUCSanRace(plugin::GenericDeviceTy &Device, SanitizerTables &Tables,
+ const __tsan_gpu_race &Race);
+
+} // namespace llvm::omp::target
+
+#endif // OFFLOAD_PLUGINS_NEXTGEN_COMMON_SANITIZER_H
diff --git a/offload/plugins-nextgen/common/include/Utils/ELF.h b/offload/plugins-nextgen/common/include/Utils/ELF.h
index dcfdb5bd7b035..0083913a45fcb 100644
--- a/offload/plugins-nextgen/common/include/Utils/ELF.h
+++ b/offload/plugins-nextgen/common/include/Utils/ELF.h
@@ -13,9 +13,18 @@
#ifndef LLVM_OPENMP_LIBOMPTARGET_PLUGINS_ELF_UTILS_H
#define LLVM_OPENMP_LIBOMPTARGET_PLUGINS_ELF_UTILS_H
+#include "llvm/ADT/SmallVector.h"
#include "llvm/Object/ELF.h"
#include "llvm/Object/ELFObjectFile.h"
+#include <cstdint>
+#include <optional>
+#include <string>
+
+namespace llvm {
+class DWARFContext;
+} // namespace llvm
+
namespace utils {
namespace elf {
@@ -39,6 +48,30 @@ getSymbolAddress(const llvm::object::ELFSymbolRef &Symbol);
llvm::Expected<std::optional<llvm::object::ELFSymbolRef>>
getSymbol(const llvm::object::ObjectFile &ELFObj, llvm::StringRef Name);
+/// A resolved source location for a single DWARF frame.
+struct SourceLocation {
+ std::string FunctionName;
+ std::string FileName;
+ uint32_t Line = 0;
+ uint32_t Column = 0;
+};
+
+/// Resolves the code address \p Addr against \p DICtx into its source frames.
+/// Returns an empty vector if no symbols are present.
+llvm::SmallVector<SourceLocation> symbolize(llvm::DWARFContext &DICtx,
+ uint64_t Addr);
+
+/// Returns the data symbol covering \p Addr in the ELF object, or an empty
+/// string if none is found.
+llvm::StringRef findDataSymbol(const llvm::object::ObjectFile &ELFObj,
+ uint64_t Addr);
+
+/// Returns the function symbol covering \p Addr in the ELF object, or an empty
+/// string if none is found. Unlike DWARF this only needs the symbol table, so
+/// it still resolves when the image lacks line tables.
+llvm::StringRef findFunctionSymbol(const llvm::object::ObjectFile &ELFObj,
+ uint64_t Addr);
+
} // namespace elf
} // namespace utils
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 7c03c916058fb..7b1dc7e324d84 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -8,11 +8,14 @@
#include "RPC.h"
+#include "Sanitizer.h"
#include "Shared/Debug.h"
#include "Shared/RPCOpcodes.h"
#include "PluginInterface.h"
+#include "sanitizer/gpu_sanitizer.h"
+
#include "shared/rpc.h"
#include "shared/rpc_opcodes.h"
#include "shared/rpc_server.h"
@@ -63,6 +66,14 @@ rpc::RPCStatus handleOffloadOpcodes(plugin::GenericDeviceTy &Device,
});
break;
}
+ case TSAN_GPU_REPORT_OPCODE: {
+ static_assert(sizeof(__tsan_gpu_race) <= sizeof(rpc::Buffer));
+ Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
+ reportGPUCSanRace(Device, Device.getRPCServer()->getSanitizerTables(),
+ *reinterpret_cast<__tsan_gpu_race *>(Buffer->data));
+ });
+ break;
+ }
default:
return rpc::RPC_UNHANDLED_OPCODE;
break;
@@ -180,13 +191,16 @@ void RPCServerTy::ServerThread::run() {
}
RPCServerTy::RPCServerTy(plugin::GenericPluginTy &Plugin)
- : Buffers(std::make_unique<void *[]>(Plugin.getNumDevices())),
+ : Sanitizers(std::make_unique<SanitizerTables>()),
+ Buffers(std::make_unique<void *[]>(Plugin.getNumDevices())),
Devices(std::make_unique<plugin::GenericDeviceTy *[]>(
Plugin.getNumDevices())),
Thread(new ServerThread(Buffers.get(), Devices.get(),
Plugin.getNumDevices(), BufferMutex, Callbacks)) {
}
+RPCServerTy::~RPCServerTy() = default;
+
llvm::Error RPCServerTy::startThread() {
Thread->startThread();
return Error::success();
diff --git a/offload/plugins-nextgen/common/src/Sanitizer.cpp b/offload/plugins-nextgen/common/src/Sanitizer.cpp
new file mode 100644
index 0000000000000..c292031e5e6b6
--- /dev/null
+++ b/offload/plugins-nextgen/common/src/Sanitizer.cpp
@@ -0,0 +1,191 @@
+//===-- Sanitizer.cpp - Host-side GPU sanitizer reporting -----------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "Sanitizer.h"
+
+#include "PluginInterface.h"
+#include "Utils/ELF.h"
+
+#include "llvm/DebugInfo/DWARF/DWARFContext.h"
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/Object/ObjectFile.h"
+#include "llvm/Support/FormatVariadic.h"
+#include "llvm/Support/raw_ostream.h"
+
+#include <optional>
+#include <utility>
+
+using namespace llvm;
+using namespace omp;
+using namespace target;
+
+namespace {
+
+// The sentinel image we use to identify the proper PC location.
+constexpr StringRef RPCClientSymbol = "__llvm_rpc_client";
+
+// Link-time ELF value of \p Name in \p Obj, if the image defines it.
+std::optional<uint64_t> getSymbolValue(const object::ObjectFile &Obj,
+ StringRef Name) {
+ Expected<std::optional<object::ELFSymbolRef>> SymOrErr =
+ utils::elf::getSymbol(Obj, Name);
+ if (!SymOrErr)
+ return consumeError(SymOrErr.takeError()), std::nullopt;
+ if (!*SymOrErr)
+ return std::nullopt;
+ Expected<uint64_t> ValueOrErr = (*SymOrErr)->getValue();
+ if (!ValueOrErr)
+ return consumeError(ValueOrErr.takeError()), std::nullopt;
+ return *ValueOrErr;
+}
+
+// A loaded image selected to symbolize a report against, holding the data
+// needed to rebase raw device addresses into image VAs.
+struct ResolvedImage {
+ std::unique_ptr<object::ObjectFile> Obj;
+ std::unique_ptr<DWARFContext> DICtx;
+ int64_t Bias = 0;
+
+ explicit operator bool() const { return Obj != nullptr; }
+ uintptr_t getAddress(uintptr_t DeviceAddr) const { return DeviceAddr + Bias; }
+};
+
+// Returns the difference between the RPC client symbol on the device and in the
+// image to get the relative offset.
+std::optional<int64_t> getImageOffset(plugin::GenericDeviceTy &Device,
+ plugin::DeviceImageTy &Image,
+ const object::ObjectFile &Obj) {
+ std::optional<uintptr_t> ImageVA = getSymbolValue(Obj, RPCClientSymbol);
+ if (!ImageVA)
+ return std::nullopt;
+ plugin::GlobalTy Client(RPCClientSymbol.str());
+ if (Error Err = Device.Plugin.getGlobalHandler().getGlobalMetadataFromDevice(
+ Device, Image, Client))
+ return consumeError(std::move(Err)), std::nullopt;
+ return static_cast<int64_t>(*ImageVA) -
+ static_cast<int64_t>(reinterpret_cast<uintptr_t>(Client.getPtr()));
+}
+
+// Identify which of the loaded images contains our program-counter. We use the
+// address of the RPC client as our sentinel to find the relative offset.
+ResolvedImage resolve(plugin::GenericDeviceTy &Device, uintptr_t PC) {
+ for (plugin::DeviceImageTy *Image : Device.LoadedImages) {
+ if (!Image)
+ continue;
+
+ Expected<std::unique_ptr<object::ObjectFile>> ObjOrErr =
+ object::ObjectFile::createObjectFile(Image->getMemoryBuffer());
+ if (!ObjOrErr) {
+ consumeError(ObjOrErr.takeError());
+ continue;
+ }
+
+ std::optional<int64_t> Bias = getImageOffset(Device, *Image, **ObjOrErr);
+ if (!Bias)
+ continue;
+
+ // Step back into the call site from the saved return address.
+ uintptr_t LookupPC = PC ? PC - 1 : PC;
+ if (utils::elf::findFunctionSymbol(**ObjOrErr, LookupPC + *Bias).empty())
+ continue;
+
+ ResolvedImage RI;
+ RI.Obj = std::move(*ObjOrErr);
+ RI.DICtx = DWARFContext::create(*RI.Obj);
+ RI.Bias = *Bias;
+ return RI;
+ }
+ return {};
+}
+
+// Print the source frames for a device PC, one line per inlined frame.
+void printBacktrace(raw_ostream &OS, const ResolvedImage &RI, uintptr_t PC) {
+ uintptr_t ImagePC = RI.getAddress(PC);
+ // The device records return addresses; step back into the call site.
+ uintptr_t LookupPC = ImagePC ? ImagePC - 1 : ImagePC;
+ SmallVector<utils::elf::SourceLocation> Frames =
+ RI ? utils::elf::symbolize(*RI.DICtx, LookupPC)
+ : SmallVector<utils::elf::SourceLocation>();
+ if (Frames.empty()) {
+ StringRef Fn =
+ RI ? utils::elf::findFunctionSymbol(*RI.Obj, LookupPC) : StringRef();
+ OS << formatv("==CSAN== #0 {0} ({1:x})\n",
+ Fn.empty() ? "??" : demangle(Fn), ImagePC);
+ return;
+ }
+ for (auto [I, Frame] : enumerate(Frames))
+ OS << formatv("==CSAN== #{0} {1} {2}:{3}:{4} ({5:x})\n", I,
+ Frame.FunctionName.empty() ? "??" : Frame.FunctionName,
+ Frame.FileName.empty() ? "??" : Frame.FileName, Frame.Line,
+ Frame.Column, ImagePC);
+}
+
+void reportRace(raw_ostream &OS, const ResolvedImage &RI,
+ const __tsan_gpu_race &Race) {
+ StringRef Op = (Race.access_type & TSAN_GPU_ACCESS_COMPOUND)
+ ? "Read-modify-write"
+ : (Race.access_type & TSAN_GPU_ACCESS_WRITE) ? "Write"
+ : "Read";
+ StringRef Atomic =
+ (Race.access_type & TSAN_GPU_ACCESS_ATOMIC) ? "atomic " : "";
+ StringRef Kind = Race.kind == TSAN_GPU_UNKNOWN_ORIGIN
+ ? "data race (unknown origin)"
+ : Race.kind == TSAN_GPU_INTRA_WAVE ? "data race (intra-wave)"
+ : "data race";
+
+ OS << formatv("==CSAN== WARNING: ConcurrencySanitizer: {0}\n", Kind);
+ OS << formatv("==CSAN== {0}{1} of size {2} at {3:x} in block "
+ "({4},{5},{6}) thread ({7},{8},{9}) lane {10}\n",
+ Atomic, Op, Race.size, Race.addr, Race.block[0], Race.block[1],
+ Race.block[2], Race.thread[0], Race.thread[1], Race.thread[2],
+ Race.lane);
+
+ printBacktrace(OS, RI, Race.pc);
+ if (Race.kind == TSAN_GPU_INTRA_WAVE) {
+ // The conflicting lane executed the same instruction, so it shares the PC.
+ OS << formatv("==CSAN== Previous access in block ({0},{1},{2}) thread "
+ "({3},{4},{5}) lane {6}\n",
+ Race.block[0], Race.block[1], Race.block[2],
+ Race.peer_thread[0], Race.peer_thread[1], Race.peer_thread[2],
+ Race.peer_lane);
+ printBacktrace(OS, RI, Race.pc);
+ } else if (Race.peer_pc) {
+ OS << "==CSAN== Previous access:\n";
+ printBacktrace(OS, RI, Race.peer_pc);
+ }
+ if (RI) {
+ StringRef Var =
+ utils::elf::findDataSymbol(*RI.Obj, RI.getAddress(Race.addr));
+ if (!Var.empty())
+ OS << formatv("==CSAN== Address {0:x} is global variable '{1}'\n",
+ Race.addr, Var);
+ }
+}
+
+} // namespace
+
+// Deduplicates races on the conflicting program counters, similar to TSan.
+bool llvm::omp::target::SanitizerTables::isNewRace(uintptr_t PC,
+ uintptr_t PeerPC,
+ unsigned Kind) {
+ if (PC > PeerPC)
+ std::swap(PC, PeerPC);
+ std::lock_guard<std::mutex> Guard(Mtx);
+ return Races.insert({PC ^ (static_cast<uint64_t>(Kind) << 56), PeerPC})
+ .second;
+}
+
+void llvm::omp::target::reportGPUCSanRace(plugin::GenericDeviceTy &Device,
+ SanitizerTables &Tables,
+ const __tsan_gpu_race &Race) {
+ // Drop duplicates before the expensive ELF/DWARF parsing below.
+ if (!Tables.isNewRace(Race.pc, Race.peer_pc, Race.kind))
+ return;
+
+ reportRace(errs(), resolve(Device, Race.pc), Race);
+}
diff --git a/offload/plugins-nextgen/common/src/Utils/ELF.cpp b/offload/plugins-nextgen/common/src/Utils/ELF.cpp
index 6b7760f3a1215..b411dca539c8d 100644
--- a/offload/plugins-nextgen/common/src/Utils/ELF.cpp
+++ b/offload/plugins-nextgen/common/src/Utils/ELF.cpp
@@ -13,6 +13,9 @@
#include "Utils/ELF.h"
#include "llvm/BinaryFormat/Magic.h"
+#include "llvm/DebugInfo/DIContext.h"
+#include "llvm/DebugInfo/DWARF/DWARFContext.h"
+#include "llvm/Demangle/Demangle.h"
#include "llvm/Object/Binary.h"
#include "llvm/Object/ELFObjectFile.h"
#include "llvm/Object/ELFTypes.h"
@@ -363,3 +366,55 @@ utils::elf::getSymbolAddress(const ELFSymbolRef &SymRef) {
return getSymbolAddressImpl(*ELFObj, SymRef);
return createError("Only 64-bit ELF files are supported");
}
+
+SmallVector<utils::elf::SourceLocation>
+utils::elf::symbolize(DWARFContext &DICtx, uint64_t Addr) {
+ SmallVector<SourceLocation> Locations;
+ DILineInfoSpecifier Spec(DILineInfoSpecifier::FileLineInfoKind::RawValue,
+ DILineInfoSpecifier::FunctionNameKind::LinkageName);
+ DIInliningInfo Info =
+ DICtx.getInliningInfoForAddress(object::SectionedAddress{Addr}, Spec);
+ Locations.reserve(Info.getNumberOfFrames());
+ for (uint32_t I = 0, N = Info.getNumberOfFrames(); I < N; ++I) {
+ const DILineInfo &Frame = Info.getFrame(I);
+ Locations.push_back({demangle(Frame.FunctionName), Frame.FileName,
+ static_cast<uint32_t>(Frame.Line),
+ static_cast<uint32_t>(Frame.Column)});
+ }
+ return Locations;
+}
+
+// Returns the name of the symbol of type \p Type whose [value, value + size)
+// range covers \p Addr, or an empty string if none is found.
+static StringRef findSymbolByAddress(const ObjectFile &ELFObj, uint64_t Addr,
+ uint8_t Type) {
+ for (const SymbolRef &Sym : ELFObj.symbols()) {
+ ELFSymbolRef ESym(Sym);
+ if (ESym.getELFType() != Type || ESym.getSize() == 0)
+ continue;
+ Expected<uint64_t> ValueOrErr = Sym.getValue();
+ if (!ValueOrErr) {
+ consumeError(ValueOrErr.takeError());
+ continue;
+ }
+ if (Addr < *ValueOrErr || Addr >= *ValueOrErr + ESym.getSize())
+ continue;
+ Expected<StringRef> NameOrErr = Sym.getName();
+ if (!NameOrErr) {
+ consumeError(NameOrErr.takeError());
+ continue;
+ }
+ if (!NameOrErr->empty())
+ return *NameOrErr;
+ }
+ return StringRef();
+}
+
+StringRef utils::elf::findDataSymbol(const ObjectFile &ELFObj, uint64_t Addr) {
+ return findSymbolByAddress(ELFObj, Addr, ELF::STT_OBJECT);
+}
+
+StringRef utils::elf::findFunctionSymbol(const ObjectFile &ELFObj,
+ uint64_t Addr) {
+ return findSymbolByAddress(ELFObj, Addr, ELF::STT_FUNC);
+}
>From 7b4f8c52851afa366bcfee4f0e806a85c246c805 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 14 Jul 2026 20:07:30 -0500
Subject: [PATCH 2/2] comments
---
compiler-rt/include/CMakeLists.txt | 2 +-
.../include/sanitizer/csan_interface.h | 47 +++++++++++++++
compiler-rt/include/sanitizer/gpu_sanitizer.h | 59 -------------------
compiler-rt/lib/csan/csan_gpu.cpp | 2 +-
.../common/include/Sanitizer.h | 2 +-
offload/plugins-nextgen/common/src/RPC.cpp | 2 +-
6 files changed, 51 insertions(+), 63 deletions(-)
create mode 100644 compiler-rt/include/sanitizer/csan_interface.h
delete mode 100644 compiler-rt/include/sanitizer/gpu_sanitizer.h
diff --git a/compiler-rt/include/CMakeLists.txt b/compiler-rt/include/CMakeLists.txt
index 05c78eeace345..dd3372e540b69 100644
--- a/compiler-rt/include/CMakeLists.txt
+++ b/compiler-rt/include/CMakeLists.txt
@@ -4,8 +4,8 @@ if (COMPILER_RT_BUILD_SANITIZERS)
sanitizer/asan_interface.h
sanitizer/common_interface_defs.h
sanitizer/coverage_interface.h
+ sanitizer/csan_interface.h
sanitizer/dfsan_interface.h
- sanitizer/gpu_sanitizer.h
sanitizer/hwasan_interface.h
sanitizer/linux_syscall_hooks.h
sanitizer/lsan_interface.h
diff --git a/compiler-rt/include/sanitizer/csan_interface.h b/compiler-rt/include/sanitizer/csan_interface.h
new file mode 100644
index 0000000000000..fc3cf4bb5ed22
--- /dev/null
+++ b/compiler-rt/include/sanitizer/csan_interface.h
@@ -0,0 +1,47 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Shared ABI contract for GPU sanitizer reports shipped to the host over RPC.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef SANITIZER_CSAN_INTERFACE_H
+#define SANITIZER_CSAN_INTERFACE_H
+
+#include <sanitizer/common_interface_defs.h>
+
+// RPC opcode identifying a GPU race report. The high byte tags the sanitizer
+// family ('s') so distinct GPU tools can share the RPC channel.
+static const unsigned TSAN_GPU_REPORT_OPCODE = ('s' << 24) | 1;
+
+// Access classification flags, mirroring the Linux kernel's KCSAN model.
+static const unsigned TSAN_GPU_ACCESS_WRITE = 1 << 0; // Non-atomic write.
+static const unsigned TSAN_GPU_ACCESS_COMPOUND = 1 << 1; // Read-modify-write.
+static const unsigned TSAN_GPU_ACCESS_ATOMIC = 1 << 2; // Atomic access.
+
+// Race report kinds.
+static const unsigned TSAN_GPU_DATA_RACE = 0; // Conflicting access.
+static const unsigned TSAN_GPU_UNKNOWN_ORIGIN = 1; // Value change.
+static const unsigned TSAN_GPU_INTRA_WAVE = 2; // Race within a wavefront.
+
+// The data associated with a single detected race.
+struct __tsan_gpu_race {
+ uint64_t pc; // Device PC of the reporting access.
+ uint64_t peer_pc; // Device PC of the conflicting access.
+ uint64_t addr; // Accessed device address.
+ uint32_t size; // Access size in bytes.
+ uint32_t access_type; // Bitwise-or of TSAN_GPU_ACCESS_* flags.
+ uint32_t kind; // One of the TSAN_GPU_* race kinds.
+ uint32_t block[3]; // Block / workgroup id.
+ uint16_t thread[3]; // Thread / work-item id within the block.
+ uint8_t lane; // Lane id within the wave.
+ uint8_t peer_lane; // Conflicting lane for intra-wave races.
+ uint16_t peer_thread[3]; // Conflicting thread id for intra-wave races.
+};
+
+#endif // SANITIZER_CSAN_INTERFACE_H
diff --git a/compiler-rt/include/sanitizer/gpu_sanitizer.h b/compiler-rt/include/sanitizer/gpu_sanitizer.h
deleted file mode 100644
index e38c677953c62..0000000000000
--- a/compiler-rt/include/sanitizer/gpu_sanitizer.h
+++ /dev/null
@@ -1,59 +0,0 @@
-//===-- sanitizer/gpu_sanitizer.h -------------------------------*- C++ -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// Shared ABI contract for GPU sanitizer reports shipped to the host over RPC.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef SANITIZER_GPU_SANITIZER_H
-#define SANITIZER_GPU_SANITIZER_H
-
-#include <stdint.h>
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-#define SANITIZER_GPU_OPCODE(n) (('s' << 24) | (n))
-
-#define TSAN_GPU_REPORT_OPCODE SANITIZER_GPU_OPCODE(1)
-
-/// Access classification flags, mirroring the Linux kernel's KCSAN model.
-enum {
- TSAN_GPU_ACCESS_WRITE = 1 << 0, ///< Non-atomic write.
- TSAN_GPU_ACCESS_COMPOUND = 1 << 1, ///< Read-modify-write.
- TSAN_GPU_ACCESS_ATOMIC = 1 << 2, ///< Atomic access.
-};
-
-/// Race report kinds.
-enum {
- TSAN_GPU_DATA_RACE = 0, ///< Conflicting access.
- TSAN_GPU_UNKNOWN_ORIGIN = 1, ///< Watched value changed with no finder.
- TSAN_GPU_INTRA_WAVE = 2, ///< Conflict between lanes of the same wave.
-};
-
-/// The data associated with a single detected race.
-typedef struct __tsan_gpu_race {
- uint64_t pc; ///< Device PC of the reporting access.
- uint64_t peer_pc; ///< Device PC of the conflicting access.
- uint64_t addr; ///< Accessed device address.
- uint32_t size; ///< Access size in bytes.
- uint32_t access_type; ///< Bitwise-or of TSAN_GPU_ACCESS_* flags.
- uint32_t kind; ///< One of the TSAN_GPU_* race kinds.
- uint32_t block[3]; ///< Block / workgroup id.
- uint16_t thread[3]; ///< Thread / work-item id within the block.
- uint8_t lane; ///< Lane id within the wave.
- uint8_t peer_lane; ///< Conflicting lane for intra-wave races.
- uint16_t peer_thread[3]; ///< Conflicting thread id for intra-wave races.
-} __tsan_gpu_race;
-
-#ifdef __cplusplus
-} // extern "C"
-#endif
-
-#endif // SANITIZER_GPU_SANITIZER_H
diff --git a/compiler-rt/lib/csan/csan_gpu.cpp b/compiler-rt/lib/csan/csan_gpu.cpp
index 27056be71fa3a..2de025ec9ba0c 100644
--- a/compiler-rt/lib/csan/csan_gpu.cpp
+++ b/compiler-rt/lib/csan/csan_gpu.cpp
@@ -14,7 +14,7 @@
#include <gpuintrin.h>
#include <stdint.h>
-#include "sanitizer/gpu_sanitizer.h"
+#include "sanitizer/csan_interface.h"
#include "sanitizer_common/sanitizer_internal_defs.h"
#include "shared/rpc.h"
diff --git a/offload/plugins-nextgen/common/include/Sanitizer.h b/offload/plugins-nextgen/common/include/Sanitizer.h
index fd6020afec567..45f0c13cdae82 100644
--- a/offload/plugins-nextgen/common/include/Sanitizer.h
+++ b/offload/plugins-nextgen/common/include/Sanitizer.h
@@ -9,7 +9,7 @@
#ifndef OFFLOAD_PLUGINS_NEXTGEN_COMMON_SANITIZER_H
#define OFFLOAD_PLUGINS_NEXTGEN_COMMON_SANITIZER_H
-#include "sanitizer/gpu_sanitizer.h"
+#include "sanitizer/csan_interface.h"
#include "llvm/ADT/DenseSet.h"
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 7b1dc7e324d84..55adcd8cda0b6 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -14,7 +14,7 @@
#include "PluginInterface.h"
-#include "sanitizer/gpu_sanitizer.h"
+#include "sanitizer/csan_interface.h"
#include "shared/rpc.h"
#include "shared/rpc_opcodes.h"
More information about the llvm-commits
mailing list