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
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* smiley.cu - vectorAdd draws a smiley face.
*
* Four vectorAdd passes execute on the GPU; each computes one scalar field
*
* c[i] = a[i] + b[i]
*
* over the same n = 32768 elements laid out as a 256x128 sample grid:
*
* pass 0 a=u^2 b=v^2 -> face disk field
* pass 1 a=(u+xe)^2 b=(v+ye)^2 -> left-eye disk field
* pass 2 a=(u-xe)^2 b=(v+ye)^2 -> right-eye disk field
* pass 3 a=u^2 b=(v-ym)^2 -> mouth annulus field
*
* with u,v centre-relative, aspect-corrected coordinates. The host copies
* the four fields back and thresholds them to render a 64x32 ASCII smiley
* (4x4 samples per character cell). The artwork therefore comes out of GPU
* addition results only - no GPU-side drawing, no CPU-side math on pixels
* beyond the thresholds.
*
* Output lines are stable and greppable (matches tools/cuda-ab/ab.sh):
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* MALLOC bytes=<n> rc=<0|err>
* VADD passes=4 n=<count> bad=<mismatches> checksum=<sum>
* RESULT cells=<W>x<H> inked=<cells> 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 smiley smiley.cu
*
* Then, from the guest: /tmp/smiley
*/
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cuda_runtime.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];
}
/* sample grid */
#define GW 256
#define GH 128
#define GN (GW * GH)
/* character cell = 4x4 samples */
#define CW 64
#define CH 32
#define CELL 4
/* feature geometry, in aspect-corrected grid units (u ~ +/-127, v ~ +/-127) */
#define FACE_R2 10000.0f /* face disk: u^2+v^2 <= 100^2 */
#define EYE_X 36.0f
#define EYE_Y 30.0f
#define EYE_R2 144.0f /* eyes: r=12 */
#define MOUTH_Y 24.0f
#define MOUTH_IN 2025.0f /* 45^2 */
#define MOUTH_OUT 3025.0f /* 55^2 */
#define MOUTH_V 40.0f
int main(void)
{
cudaDeviceProp prop;
float *h_a, *h_b, *h_out, *fields;
float *d_a, *d_b, *d_c;
size_t bytes = (size_t)GN * sizeof(float);
double checksum = 0.0;
int bad = 0, pass, i;
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);
{
cudaError_t e_a = cudaMalloc(&d_a, bytes);
cudaError_t e_b = e_a == cudaSuccess ?
cudaMalloc(&d_b, bytes) : e_a;
cudaError_t e_c = e_b == cudaSuccess ?
cudaMalloc(&d_c, bytes) : e_b;
printf("MALLOC bytes=%zu rc=%d\n", bytes, (int)e_a);
if (e_a != cudaSuccess) {
size_t fm = 0, tm = 0;
cudaMemGetInfo(&fm, &tm);
printf("MEMINFO free=%llu total=%llu\n",
(unsigned long long)fm,
(unsigned long long)tm);
}
CHECK(e_a);
CHECK(e_b);
CHECK(e_c);
}
h_a = (float *)malloc(bytes);
h_b = (float *)malloc(bytes);
h_out = (float *)malloc(bytes);
fields = (float *)malloc((size_t)4 * bytes);
if (!h_a || !h_b || !h_out || !fields)
return 1;
for (pass = 0; pass < 4; pass++) {
for (i = 0; i < GN; i++) {
int x = i % GW, y = i / GW;
float u = (float)x - 127.5f;
float v = ((float)y - 63.5f) * 2.0f;
float pa, pb;
switch (pass) {
case 0: /* face */
pa = u * u; pb = v * v;
break;
case 1: /* left eye */
pa = (u + EYE_X) * (u + EYE_X);
pb = (v + EYE_Y) * (v + EYE_Y);
break;
case 2: /* right eye */
pa = (u - EYE_X) * (u - EYE_X);
pb = (v + EYE_Y) * (v + EYE_Y);
break;
default: /* mouth */
pa = u * u;
pb = (v - MOUTH_Y) * (v - MOUTH_Y);
break;
}
h_a[i] = pa;
h_b[i] = pb;
}
CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
vadd<<<(GN + 255) / 256, 256>>>(d_a, d_b, d_c, GN);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_out, d_c, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < GN; i++) {
if (h_out[i] != h_a[i] + h_b[i])
bad++;
checksum += (double)h_out[i];
}
for (i = 0; i < GN; i++)
fields[(size_t)pass * GN + i] = h_out[i];
}
printf("VADD passes=4 n=%d bad=%d checksum=%.0f\n", GN, bad, checksum);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
{
static char art[CH][CW + 1];
int inked = 0, cy, cx;
for (cy = 0; cy < CH; cy++) {
for (cx = 0; cx < CW; cx++) {
int hit = 0, sy, sx;
const float *f0 = fields;
const float *f1 = fields + GN;
const float *f2 = fields + 2 * GN;
const float *f3 = fields + 3 * GN;
for (sy = 0; sy < CELL; sy++) {
int y = cy * CELL + sy;
float v = ((float)y - 63.5f) * 2.0f;
for (sx = 0; sx < CELL; sx++) {
int x = cx * CELL + sx;
int idx = y * GW + x;
int inside;
inside = f0[idx] <= FACE_R2;
if (f1[idx] <= EYE_R2 ||
f2[idx] <= EYE_R2)
inside = 0;
if (f3[idx] >= MOUTH_IN &&
f3[idx] <= MOUTH_OUT &&
v >= MOUTH_V)
inside = 0;
hit += inside;
}
}
art[cy][cx] = hit >= CELL * CELL / 2 ? '#' :
' ';
inked += art[cy][cx] == '#';
}
art[cy][CW] = 0;
}
printf("RESULT cells=%dx%d inked=%d checksum=%.0f\n",
CW, CH, inked, checksum);
for (cy = 0; cy < CH; cy++)
printf("%s\n", art[cy]);
}
printf("STATUS OK\n");
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
free(h_a);
free(h_b);
free(h_out);
free(fields);
return 0;
}
|
