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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* p2_managed_probe - EXPECTED-FAILURE PROBE (documentation, not acceptance).
*
* Managed (UVM) memory is a known unsupported path for the emulated vGPU.
* This probe attempts cudaMallocManaged, records the outcome as a PROBE
* line, clears any error, and then proves that ordinary cudaMalloc + launch
* still works afterwards. The emulated guest is expected to report a
* failure here; the host oracle may report success. p2/run_guest.sh treats
* this program as informational (content not diffed), but requires the
* program to complete with STATUS OK - i.e. cleanly report the managed
* behaviour and keep working - so a hang, crash or poisoned context still
* fails the suite.
*
* Device memory: 4 KB managed (if it works) + 4 KB normal.
*/
#include "common.h"
__global__ void fillk(int *p, int v, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
p[i] = v + i;
}
int main(int argc, char **argv)
{
if (p2_emit_mode(argc, argv)) {
printf("PROBE managed state=no-oracle-gpu note=emit-mode\n");
printf("STATUS OK\n");
return 0;
}
p2_device_info();
int *p = NULL;
cudaError_t e = cudaMallocManaged(&p, sizeof(int) * 4);
if (e == cudaSuccess) {
cudaError_t e2, e3;
int ok = 0;
fillk<<<1, 1>>>(p, 5, 4);
e2 = cudaGetLastError();
e3 = cudaDeviceSynchronize();
if (e2 == cudaSuccess && e3 == cudaSuccess)
ok = (p[0] == 5 && p[1] == 6 && p[2] == 7 &&
p[3] == 8) ? 1 : 0;
printf("PROBE managed alloc=ok write_read=%s\n",
ok ? "ok" : "bad");
cudaFree(p);
} else {
printf("PROBE managed alloc=fail rc=%d name=%s\n", (int)e,
cudaGetErrorString(e));
(void)cudaGetLastError();
}
int *d = NULL, bad = 0;
cudaError_t e4 = cudaMalloc(&d, sizeof(int) * 4);
if (e4 == cudaSuccess) {
int h[4];
fillk<<<1, 1>>>(d, 100, 4);
(void)cudaGetLastError();
cudaDeviceSynchronize();
cudaMemcpy(h, d, sizeof(h), cudaMemcpyDeviceToHost);
for (int i = 0; i < 4; i++)
if (h[i] != 100 + i)
bad++;
cudaFree(d);
printf("PROBE fallback rc=0 bad=%d\n", bad);
} else {
printf("PROBE fallback rc=%d name=%s\n", (int)e4,
cudaGetErrorString(e4));
bad = 1;
}
return p2_status(bad);
}
|
