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
|
/* Copyright 2024 Joshua Bakita
* Simple kernel that clocks how long a kernel launch takes to the GPU.
* Prints samples to stdout and summary statistics to stderr.
*/
#include <time.h>
#include <stdio.h>
#include <stdint.h>
#include <unistd.h>
#include "testbench.h"
__global__ void flag_on_gpu(volatile int *flag) {
*flag = 1;
}
// Get launch overhead
long measure_launch_overhead() {
volatile int *barrier;
struct timespec start, end;
SAFE(cudaHostAlloc(&barrier, sizeof(*barrier), cudaHostAllocMapped));
*barrier = 0;
clock_gettime(CLOCK_MONOTONIC_RAW, &start);
flag_on_gpu<<<100,100>>>(barrier);
while (!*barrier) continue;
clock_gettime(CLOCK_MONOTONIC_RAW, &end);
SAFE(cudaDeviceSynchronize());
return timediff(start, end);
}
int main(int argc, char **argv) {
int *__unused, i;
unsigned long num_iters;
long time;
long double cumulative_time = 0;
if (argc != 2 || !strcmp(argv[1], "--help") || !strcmp(argv[1], "-h")) {
fprintf(stderr, "Usage: %s [# of samples]\n", argv[0]);
return 1;
}
num_iters = strtoul(argv[1], NULL, 10);
// Initialize CUDA and a context (hack)
SAFE(cudaMalloc(&__unused, 8));
// Run once to ensure the kernel is compiled
time = measure_launch_overhead();
fprintf(stderr, "(%d) %ld ns (%.2f ms) warmup launch overhead\n", getpid(),
time, ns2ms((double)time));
// Time several kernel launches
for (i = 0; i < num_iters; i++) {
time = measure_launch_overhead();
cumulative_time += time;
// Print one sample per line
fprintf(stdout, "%ld\n", time);
}
fprintf(stderr, "(%d) %.0Lf ns (%.2Lf ms) average launch overhead\n", getpid(),
cumulative_time / num_iters, ns2ms(cumulative_time / num_iters));
return 0;
}
|
