================ @@ -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) ---------------- yxsamliu wrote:
`size` can be `UINT32_MAX`, while `i` is `u32`. Near the end, `i + 8` can wrap, so this loop may never finish and may read out of bounds. Could we use a wider index or an overflow-safe condition and add a focused test? https://github.com/llvm/llvm-project/pull/225782 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
