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
152
153
154
155
156
157
158
|
/* Copyright 2023 Joshua Bakita
* Tool that tracks preemptions as discontinuities in GPU time and logs them.
*
* The logged intervals are output to standard out as a CSV, with each row
* containing a single interval. The first column contains the interval start
* time in nanoseconds (ns), and the second column contains the interval end
* time in ns.
*
* All progress and error messages and printed to standard error.
*
* Assumptions:
* - The GPU clock continues to tick, regardless of if this is the current
* application running.
*/
#include <stdio.h>
#include <stdint.h>
#include <sys/types.h>
#include <time.h>
#include <unistd.h>
#include "task_host_utilities.cu"
#include "testbench.h"
struct interval {
uint64_t start;
uint64_t end;
};
// Minimum time discontinuity which indicates a gap between intervals
// One clock tick is about 1ns
#define MIN_PREEMPT_TICKS 2*1000 // ~2us
// Watch for discontinuities in the GPU clock, indicating intervals during
// which we we were preempted and another context ran.
// @param ivls Base address of array to log interval start and end times
// @param num_ivls How many intervals should be logged before exiting?
__global__ void loop_on_gpu(struct interval *ivls, int num_ivls) {
struct interval *curr_ivl = ivls;
uint64_t last_time = GlobalTimer64();
if (num_ivls <= 0)
return;
curr_ivl->start = last_time;
// While no preemptions are occuring, this repeatedly obtains the clock,
// storing the value in `last_time`. When we're switched away from and back
// to, there will appear to be a large gap between `now` and `last_time`.
// We record `last_time` as the end of the last interval, and `now` as the
// start of the next interval, then resume repeatedly obtaining the clock.
// Continue to log intervals until the pointer `curr_ivl` points to a
// location off the end of the `ivls` array.
while (curr_ivl < ivls + num_ivls) {
uint64_t now = GlobalTimer64();
if (now > last_time + MIN_PREEMPT_TICKS) {
curr_ivl->end = last_time;
curr_ivl++;
curr_ivl->start = now;
}
last_time = now;
}
}
static const char* usage_msg = "\
Usage: %s NUM_INTERVALS [-r]\n\
Spin on the GPU, logging intervals to stdout during which we are scheduled.\n\
NUM_INTERVALS Number of intervals of execution to log (logs one less than\n\
this number of preemptions).\n\
-r, --raw Print raw logged GPU times (skip conversion to CPU time).\n";
int main(int argc, char **argv) {
cudaError_t err; // Needed for the SAFE() macro
struct interval *ivls_gpu, *ivls;
struct timespec start, end, end_ivls_only;
int num_ivls, skip_conversion;
uint64_t dev_ns = 0;
double d2h_scale = 1.0, host_s = 0;
pid_t pid = getpid();
// Argument parsing (could be improved)
if (argc != 2 && argc != 3) {
fprintf(stderr, usage_msg, argv[0]);
return 1;
}
num_ivls = atoi(argv[1]);
if (!num_ivls) {
fprintf(stderr, usage_msg, argv[0]);
return 1;
}
skip_conversion = argc == 3 &&
(!strcmp(argv[2], "-r") || !strcmp(argv[2], "--raw"));
// Synchronize the GPU and CPU clocks (if requested) using the utilities
// from task_host_utilities.cu.
if (!skip_conversion) {
// Note that this appears to almost always be 1.0 on recent GPUs
d2h_scale = GetGPUTimerScale(0);
// The following uses a subprocess for clock synchronization. Sketch?
GetHostDeviceTimeOffset(0, &host_s, &dev_ns);
}
// Initialize GPU and CPU memory to store the intervals of execution
// Note that we use GPU memory rather than pinned host memory in order to make
// the time-monitoring loop on-GPU as fast as possible. (Pinned host memory is
// slower to read/write from/to.)
SAFE(cudaMalloc(&ivls_gpu, num_ivls * sizeof(struct interval)));
ivls = (struct interval*)malloc(num_ivls * sizeof(struct interval));
if (!ivls) {
perror("While allocating memory");
return 1;
}
// Monitor execution intervals
fprintf(stderr, "(%d) Started monitoring execution for %d intervals...\n",
pid, num_ivls);
clock_gettime(CLOCK_MONOTONIC_RAW, &start);
loop_on_gpu<<<1,1>>>(ivls_gpu, num_ivls);
SAFE(cudaDeviceSynchronize());
clock_gettime(CLOCK_MONOTONIC_RAW, &end_ivls_only);
// Pull back results onto the CPU
SAFE(cudaMemcpy(ivls, ivls_gpu, num_ivls * sizeof(struct interval),
cudaMemcpyDeviceToHost));
clock_gettime(CLOCK_MONOTONIC_RAW, &end);
// Print out the intervals, the start and end times of our monitor, and some
// summary statistics
for (int i = 0; i < num_ivls; i++) {
if (skip_conversion) {
printf("%lu, %lu\n", ivls[i].start, ivls[i].end);
} else {
// Convert all the interval times into CPU time.
// First compute how far this interval's starting point is, in CPU seconds,
// from the point where the device and CPU clocks were synchronized
uint64_t start_time_since_sync = (ivls[i].start - dev_ns) * d2h_scale;
// Then convert into absolute CPU time via offset adjustment
uint64_t start = s2ns(host_s) + start_time_since_sync;
// Repeat the same process for interval end times
uint64_t end_time_since_sync = (ivls[i].end - dev_ns) * d2h_scale;
uint64_t end = s2ns(host_s) + end_time_since_sync;
printf("%lu, %lu\n", start, end);
}
}
fprintf(stderr, "(%d) Total monitor runtime: %ld ms\n", pid,
ns2ms(timediff(start, end)));
fprintf(stderr, "(%d) CPU logging start: %ld ns\n", pid, time2ns(start));
fprintf(stderr, "(%d) CPU logging end: %ld ns\n", pid, time2ns(end));
fprintf(stderr, "(%d) CPU copy out end: %ld ns\n", pid, time2ns(end));
if (!skip_conversion) {
// Difference between CPU time of the start of the first interval and launch time
int64_t launch_oh = s2ns(host_s) +
(ivls[0].start - dev_ns) * d2h_scale - time2ns(start);
fprintf(stderr, "(%d) Aprox launch overhead: %ld ns\n", pid, launch_oh);
fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", pid,
(long)s2ns(host_s) - dev_ns);
fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.3f\n", pid, d2h_scale);
}
return 0;
}
|
