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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* memcpy_roundtrip.cu - 8 MB H2D + kernel + D2H roundtrip plus memset check.
*
* Two data paths are exercised over an 8 MB (2M float) device buffer:
*
* 1. host -> device copy, a kernel (out[i] = in[i] + 0.5), device -> host
* copy; the host verifies the full chain element by element and sums the
* result as a checksum.
* 2. cudaMemset of a second 8 MB buffer with a known byte pattern, copied
* back and verified byte-for-byte.
*
* Output lines are stable and greppable:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* ROUNDTRIP bytes=<n> bad=<mismatches> csum=<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 memcpy_roundtrip memcpy_roundtrip.cu
*
* Then, from the guest: /tmp/memcpy_roundtrip
*/
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cstring>
#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)
#define MemsetByte 0x5A
__global__ void bump(const float *in, float *out, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
out[i] = in[i] + 0.5f;
}
int main(void)
{
cudaDeviceProp prop;
const size_t bytes = 8u << 20; /* 8 MB */
const int n = (int)(bytes / sizeof(float));
unsigned char *h_m;
float *h_in, *h_out;
float *d_in, *d_out;
unsigned char *d_m;
double csum = 0.0;
int bad = 0, i;
const int grid = (n + 255) / 256;
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);
h_in = (float *)malloc(bytes);
h_out = (float *)malloc(bytes);
h_m = (unsigned char *)malloc(bytes);
if (!h_in || !h_out || !h_m)
return 1;
for (i = 0; i < n; i++)
h_in[i] = (float)(i % 1024) * 0.25f;
CHECK(cudaMalloc(&d_in, bytes));
CHECK(cudaMalloc(&d_out, bytes));
CHECK(cudaMalloc(&d_m, bytes));
/* H2D -> kernel -> D2H roundtrip. */
CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));
bump<<<grid, 256>>>(d_in, d_out, n);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < n; i++) {
if (h_out[i] != h_in[i] + 0.5f)
bad++;
csum += (double)h_out[i];
}
/* cudaMemset pattern, copied back and verified byte-for-byte. */
CHECK(cudaMemset(d_m, MemsetByte, bytes));
CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < (int)bytes; i++) {
if (h_m[i] != (unsigned char)MemsetByte)
bad++;
}
printf("ROUNDTRIP bytes=%zu bad=%d csum=%.6f\n", bytes, bad, csum);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
printf("STATUS OK\n");
CHECK(cudaFree(d_in));
CHECK(cudaFree(d_out));
CHECK(cudaFree(d_m));
free(h_in);
free(h_out);
free(h_m);
return 0;
}
|
