nv-compute/gpu_src/gpu_patch_block_divergence_analysis.cu at 4dc3e7454a970d9592f8a039c1badf9ce4e3f1e9 · AccelProf/nv-compute · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
#include "gpu_patch.h"

#include <sanitizer_patching.h>

#include "gpu_utils.h"
#include <cstdio>

static __device__ __inline__
uint32_t GetBufferIndex(MemoryAccessTracker* pTracker) {
    uint32_t idx = MEMORY_ACCESS_BUFFER_SIZE;

    while (idx >= MEMORY_ACCESS_BUFFER_SIZE) {
        idx = atomicAdd(&(pTracker->currentEntry), 1);

        if (idx >= MEMORY_ACCESS_BUFFER_SIZE) {
            // buffer is full, wait for last writing thread to flush
            while (*(volatile uint32_t*)&(pTracker->currentEntry) >= MEMORY_ACCESS_BUFFER_SIZE);
        }
    }

    return idx;
}

static __device__ __inline__
void IncrementNumEntries(MemoryAccessTracker* pTracker) {
    DoorBell* doorbell = pTracker->doorBell;
    __threadfence();
    const uint32_t numEntries = atomicAdd((int*)&(pTracker->numEntries), 1);

    if (numEntries == MEMORY_ACCESS_BUFFER_SIZE - 1) {
        // make sure everything is visible in memory
        __threadfence_system();
        doorbell->full = true;
        while (doorbell->full);

        pTracker->numEntries = 0;
        __threadfence();
        pTracker->currentEntry = 0;
    }
}

static __device__
SanitizerPatchResult CommonCallback(
    void* userdata,
    uint64_t pc,
    void* ptr,
    uint32_t accessSize,
    uint32_t flags,
    MemoryType type)
{
    auto* pTracker = (MemoryAccessTracker*)userdata;

    uint32_t active_mask = __activemask();
    uint32_t laneid = get_laneid();
    uint32_t first_laneid = __ffs(active_mask) - 1;

    MemoryAccess* accesses = nullptr;

    if (laneid == first_laneid) {
        uint32_t idx = GetBufferIndex(pTracker);
        accesses = &pTracker->access_buffer[idx];
        accesses->accessSize = accessSize;
        accesses->flags = flags;
        accesses->warpId = get_warpid();
        accesses->ctaId = get_ctaid_as_uint64();
        accesses->pc = pc;
        accesses->active_mask = active_mask;
        accesses->type = type;
    }

    __syncwarp(active_mask);

    accesses = (MemoryAccess*) shfl((uint64_t)accesses, first_laneid, active_mask);
    if (accesses) {
        accesses->addresses[laneid] = (uint64_t)(uintptr_t)ptr;
    }

    __syncwarp(active_mask);

    if (laneid == first_laneid) {
        IncrementNumEntries(pTracker);
    }

    return SANITIZER_PATCH_SUCCESS;
}

extern "C" __device__ __noinline__
SanitizerPatchResult MemoryGlobalAccessCallback(
    void* userdata,
    uint64_t pc,
    void* ptr,
    uint32_t accessSize,
    uint32_t flags,
    const void *pData)
{
    return CommonCallback(userdata, pc, ptr, accessSize, flags, MemoryType::Global);
}

extern "C" __device__ __noinline__
SanitizerPatchResult MemorySharedAccessCallback(
    void* userdata,
    uint64_t pc,
    void* ptr,
    uint32_t accessSize,
    uint32_t flags,
    const void *pData)
{
    return CommonCallback(userdata, pc, ptr, accessSize, flags, MemoryType::Shared);
}

extern "C" __device__ __noinline__
SanitizerPatchResult MemoryLocalAccessCallback(
    void* userdata,
    uint64_t pc,
    void* ptr,
    uint32_t accessSize,
    uint32_t flags,
    const void *pData)
{
    return CommonCallback(userdata, pc, ptr, accessSize, flags, MemoryType::Local);
}

extern "C" __device__ __noinline__
SanitizerPatchResult MemcpyAsyncCallback(void* userdata, uint64_t pc, void* src, uint32_t dst, uint32_t accessSize)
{
    if (src)
    {
        CommonCallback(userdata, pc, src, accessSize, SANITIZER_MEMORY_DEVICE_FLAG_READ, MemoryType::Global);
    }

    return CommonCallback(userdata, pc, (void*)dst, accessSize, SANITIZER_MEMORY_DEVICE_FLAG_WRITE, MemoryType::Shared);
}

extern "C" __device__ __noinline__
SanitizerPatchResult BlockExitCallback(void* userdata, uint64_t pc)
{
    MemoryAccessTracker* tracker = (MemoryAccessTracker*)userdata;
    DoorBell* doorbell = tracker->doorBell;

    uint32_t active_mask = __activemask();
    uint32_t laneid = get_laneid();
    uint32_t first_laneid = __ffs(active_mask) - 1;
    int32_t pop_count = __popc(active_mask);

    if (laneid == first_laneid) {
        atomicAdd((int*)&doorbell->num_threads, -pop_count);
    }
    __syncwarp(active_mask);

    return SANITIZER_PATCH_SUCCESS;
}