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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* wedge_bisect.cu - bisect the wedge trigger: which inter-launch activity
* makes the SECOND vadd launch no-execute?
*
* Every iteration is: mutate host inputs -> H2D da/db -> launch -> sync ->
* D2H dc -> verify (the wedge_recover shape). Modes add activity between
* the verify and the next launch:
*
* mode A: none (pure second-launch class)
* mode B: an extra H2D of the same buffers (CE write before the launch)
* mode C: an extra D2H readback (CE read before the launch)
* mode D: a 500 ms sleep (no CE traffic)
*
* A no-execute launch is reported (bad>0) not masked: the inputs mutate every
* iteration.
*
* Build: nvcc -O2 --cudart=static -o wedge_bisect wedge_bisect.cu
*/
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>
#include <unistd.h>
#define CHECK(x) \
do { \
cudaError_t e_ = (x); \
if (e_ != cudaSuccess) { \
printf("CUDA_ERROR %s at line %d\n", \
cudaGetErrorString(e_), __LINE__); \
printf("STATUS FAIL\n"); \
return 1; \
} \
} while (0)
__global__ void vadd(const float *a, const float *b, float *c, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
c[i] = a[i] + b[i];
}
int main(int argc, char **argv)
{
const int n = 1 << 15;
size_t bytes = (size_t)n * sizeof(float);
float *ha, *hb, *hc;
float *da, *db, *dc;
const char *mode = argc > 1 ? argv[1] : "A";
int iters = argc > 2 ? atoi(argv[2]) : 4;
int i, bad = 0;
setvbuf(stdout, NULL, _IONBF, 0);
ha = (float *)malloc(bytes);
hb = (float *)malloc(bytes);
hc = (float *)malloc(bytes);
if (!ha || !hb || !hc)
return 1;
for (i = 0; i < n; i++) {
ha[i] = (float)(i % 977) * 0.25f;
hb[i] = (float)(i % 251) * 0.5f;
}
CHECK(cudaMalloc(&da, bytes));
CHECK(cudaMalloc(&db, bytes));
CHECK(cudaMalloc(&dc, bytes));
CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
for (i = 0; i < iters; i++) {
int j, itbad = 0;
if (i > 0) {
ha[i % n] += 1.0f;
hb[(i * 7 + 3) % n] += 2.0f;
}
CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
for (j = 0; j < n; j++)
if (hc[j] != ha[j] + hb[j])
itbad++;
printf("MODE %s iter=%d bad=%d/%d\n", mode, i, itbad, n);
bad += itbad;
if (i + 1 < iters) {
if (mode[0] == 'B') {
CHECK(cudaMemcpy(da, ha, bytes,
cudaMemcpyHostToDevice));
} else if (mode[0] == 'C') {
CHECK(cudaMemcpy(hc, dc, bytes,
cudaMemcpyDeviceToHost));
} else if (mode[0] == 'D') {
usleep(500000);
}
}
}
printf("STATUS %s\n", bad ? "FAIL" : "OK");
return bad ? 1 : 0;
}
|
