================ @@ -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]; ---------------- yxsamliu wrote:
`block_count` counts only full workgroups. A non-zero remainder adds another partial group, so this stride can be too small. For global size `(3, 2)` and workgroup size `(2, 1)`, groups `(1, 0)` and `(0, 1)` get the same ID. Could we follow libclc’s `get_num_groups` logic and add a multidimensional partial-group 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
