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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* victory_mail.cu - the Open-IOV victory program.
*
* Run from the guest. Two kernels execute on the physical GPU:
*
* mail_kernel - the GPU itself writes every character of the message into
* device memory; the host only copies the bytes back.
* vadd - the classic vectorAdd acceptance check, verified element
* by element against the host-computed result.
*
* The host prints the message only after it has been authored by GPU
* execution and copied back byte-for-byte, so the line below is proof that a
* kernel really ran on the device:
*
* the little gpu man says you've got mail
*
* Output lines are stable and greppable:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* MAIL <message>
* VADD n=<count> bad=<mismatches> checksum=<sum>
* STATUS OK|FAIL
*
* Build with a static CUDA runtime (--cudart=static) so the guest only needs
* the driver's libcuda.so, not a CUDA toolkit installation:
*
* nvcc -O2 --cudart=static -o victory_mail victory_mail.cu
*
* Then, from the guest: /tmp/victory_mail
*/
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>
#define MAIL "the little gpu man says you've got mail"
#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)
/* The message lives in device code: the GPU both provides and writes it. */
__device__ const char kMail[] = MAIL;
/* Each thread writes exactly one character of the message. */
__global__ void mail_kernel(char *out, int len)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < len)
out[i] = kMail[i];
}
__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(void)
{
cudaDeviceProp prop;
const int len = (int)strlen(MAIL);
const int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
char *d_mail, h_mail[128];
float *h_a, *h_b, *h_c;
float *d_a, *d_b, *d_c;
double sum = 0.0;
int bad = 0, i;
/* Unbuffered: progress must be visible even if a later call blocks. */
setvbuf(stdout, NULL, _IONBF, 0);
CHECK(cudaGetDeviceProperties(&prop, 0));
printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
prop.major, prop.minor, prop.multiProcessorCount,
(unsigned long long)prop.totalGlobalMem);
/* The GPU writes the mail. */
CHECK(cudaMalloc(&d_mail, len));
mail_kernel<<<1, 64>>>(d_mail, len);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
memset(h_mail, 0, sizeof(h_mail));
CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
if (memcmp(h_mail, MAIL, len) != 0) {
printf("MAIL bad=<kernel output did not match>\n");
printf("STATUS FAIL\n");
return 1;
}
printf("MAIL %s\n", h_mail);
/* The classic acceptance check: vectorAdd. */
h_a = (float *)malloc(bytes);
h_b = (float *)malloc(bytes);
h_c = (float *)malloc(bytes);
if (!h_a || !h_b || !h_c)
return 1;
for (i = 0; i < n; i++) {
h_a[i] = (float)(i % 1000) * 0.5f;
h_b[i] = (float)(i % 997) * 0.25f;
}
CHECK(cudaMalloc(&d_a, bytes));
CHECK(cudaMalloc(&d_b, bytes));
CHECK(cudaMalloc(&d_c, bytes));
CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < n; i++) {
if (h_c[i] != h_a[i] + h_b[i])
bad++;
sum += h_c[i];
}
printf("VADD n=%d bad=%d checksum=%.6f\n", n, bad, sum);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
printf("the little gpu man says you've got mail\n");
printf("STATUS OK\n");
CHECK(cudaFree(d_mail));
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
free(h_a);
free(h_b);
free(h_c);
return 0;
}
|
