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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* p2_printf_err - device printf, a passing device assert, and an intentional
* launch error path.
*
* 1. A 4-thread kernel prints deterministic P2PRINT lines from the device.
* 2. A kernel with assert(n == 4) exercises the device-side assert
* machinery on its passing path (a fired assert is deliberately not
* tested: it traps the kernel and poisons the context).
* 3. A kernel launched with gridDim.x == 0 must be rejected:
* cudaErrorInvalidConfiguration (9), and the error must be cleared by
* cudaGetLastError. A blockDim.x == 0 launch is exercised too.
* 4. A normal launch after the errors proves error recovery.
*
* Expected error strings come from the statically linked CUDA runtime, so
* they must be identical on host and guest.
* Device memory: out 4 B.
*/
#include "common.h"
#include <cassert>
__global__ void hello(int nthreads)
{
int i = threadIdx.x;
if (i < nthreads)
printf("P2PRINT hello i=%d\n", i);
}
__global__ void second(void)
{
if (threadIdx.x == 0)
printf("P2PRINT second c=%c v=%d\n", 'Z', 7);
}
__global__ void simple(int *out)
{
out[0] = 42;
}
__global__ void assertk(int *out, int n)
{
assert(n == 4);
if (threadIdx.x == 0)
out[0] = n;
}
int main(int argc, char **argv)
{
if (p2_emit_mode(argc, argv)) {
printf("P2PRINT hello i=0\n");
printf("P2PRINT hello i=1\n");
printf("P2PRINT hello i=2\n");
printf("P2PRINT hello i=3\n");
printf("P2PRINT second c=Z v=7\n");
printf("ERR launch grid0 rc=%d name=%s\n",
(int)cudaErrorInvalidConfiguration,
cudaGetErrorString(cudaErrorInvalidConfiguration));
printf("ERR launch block0 rejected=1\n");
printf("ERR cleared rc=0\n");
printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n");
printf("RESULT assert n=4 bad=0\n");
return p2_status(0);
}
p2_device_info();
int *d_out = NULL;
P2_CHECK(cudaMalloc(&d_out, sizeof(int)));
hello<<<1, 4>>>(4);
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaDeviceSynchronize());
second<<<1, 4>>>();
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaDeviceSynchronize());
(void)cudaGetLastError();
int badgrid = 0, badblock = 0;
cudaError_t e_grid, e_block, cleared;
int got = 0;
simple<<<badgrid, 1>>>(d_out);
e_grid = cudaGetLastError();
printf("ERR launch grid0 rc=%d name=%s\n", (int)e_grid,
cudaGetErrorString(e_grid));
simple<<<1, badblock>>>(d_out);
e_block = cudaGetLastError();
printf("ERR launch block0 rejected=%d\n",
(e_block != cudaSuccess) ? 1 : 0);
cleared = cudaGetLastError();
printf("ERR cleared rc=%d\n", (int)cleared);
simple<<<1, 1>>>(d_out);
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaDeviceSynchronize());
P2_CHECK(cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost));
int grid0_bad = (e_grid != cudaErrorInvalidConfiguration) ? 1 : 0;
int block0_bad = (e_block == cudaSuccess) ? 1 : 0;
int clear_bad = (cleared != cudaSuccess) ? 1 : 0;
int recovered = (got == 42) ? 1 : 0;
int assert_bad = 0;
assertk<<<1, 1>>>(d_out, 4);
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaDeviceSynchronize());
{
int agot = 0;
P2_CHECK(cudaMemcpy(&agot, d_out, sizeof(int),
cudaMemcpyDeviceToHost));
assert_bad = (agot == 4) ? 0 : 1;
}
int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1) +
assert_bad;
printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n",
grid0_bad, block0_bad, clear_bad, recovered);
printf("RESULT assert n=4 bad=%d\n", assert_bad);
cudaFree(d_out);
return p2_status(bad);
}
|
