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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure.
*
* cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report:
* the first bad offset, and a coverage map at 4 KB granularity (correct /
* wrong / untouched-malloc-prefill) so "fill under-delivered" is
* distinguishable from "readback delivered wrong pages".
*
* Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.cu
*/
#include <cstdio>
#include <cstdlib>
#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 PAGE 4096
#define MemsetByte 0x5A
int main(void)
{
const size_t bytes = 8u << 20;
unsigned char *h;
unsigned char *d;
size_t i;
int bad = 0;
size_t first_bad = (size_t)-1;
size_t n_pages, p;
setvbuf(stdout, NULL, _IONBF, 0);
h = (unsigned char *)malloc(bytes);
if (!h)
return 1;
memset(h, 0xEE, bytes);
CHECK(cudaMalloc(&d, bytes));
printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes);
CHECK(cudaMemset(d, MemsetByte, bytes));
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < bytes; i++) {
if (h[i] != (unsigned char)MemsetByte) {
if (first_bad == (size_t)-1)
first_bad = i;
bad++;
}
}
printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n",
first_bad, first_bad, bad, bytes);
n_pages = bytes / PAGE;
for (p = 0; p < n_pages; p++) {
int good = 1, prefill = 1;
for (i = 0; i < PAGE; i++) {
if (h[p * PAGE + i] != (unsigned char)MemsetByte)
good = 0;
if (h[p * PAGE + i] != 0xEE)
prefill = 0;
if (!good && !prefill)
break;
}
if (!good) {
printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n",
p, prefill ? "PREFILL" : "OTHER",
h[p * PAGE], h[p * PAGE + 1],
h[p * PAGE + 2], h[p * PAGE + 3],
h[p * PAGE + 4095]);
if (p > 24)
break;
}
}
printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
return 0;
}
|
