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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* memset2d_probe.cu - per-phase diagnostics for victory_memset2d.
*
* victory_memset2d reports one aggregate bad count (~12.5M) that mixes all
* three phases. This probe runs the SAME three phases and reports each
* phase's bad count, the first mismatching byte (offset, got, want), and a
* per-phase 1MB-chunk histogram so the failing phase and the failing region
* are attributable in one run.
*
* Phases:
* P1 memset : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback
* P2 roundtrip: H2D 2D pattern; D2H 2D readback
* P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback
*
* Output lines:
* DEVICE ...
* PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w>
* PHIST p=1 mb=<m> ok=<n> bad=<n>
* MEMSET2D2 bad_total=<n>
* STATUS OK|FAIL
*/
#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 IMG_W 4096
#define IMG_H 1024
#define FILL_BYTE 0xA5
#define QUAD_BYTE 0x3C
static int phase_check(const char *tag, const unsigned char *back,
const unsigned char *want, size_t n, int p)
{
size_t i, first = (size_t)-1;
int bad = 0, got = 0, wantv = 0;
size_t mb;
for (i = 0; i < n; i++) {
if (back[i] != want[i]) {
if (first == (size_t)-1) {
first = i;
got = back[i];
wantv = want[i];
}
bad++;
}
}
printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n",
p, bad, first, got, wantv);
if (bad) {
size_t nb = (n + (1u << 20) - 1) >> 20;
for (mb = 0; mb < nb && mb < 8; mb++) {
size_t lo = mb << 20, hi = lo + (1u << 20);
size_t j;
int bok = 0, bbad = 0;
if (hi > n)
hi = n;
for (j = lo; j < hi; j++) {
if (back[j] == want[j])
bok++;
else
bbad++;
}
printf("PHIST p=%d mb=%zu ok=%d bad=%d\n",
p, mb, bok, bbad);
}
}
(void)tag;
return bad;
}
int main(void)
{
cudaDeviceProp prop;
const size_t w = IMG_W, h = IMG_H;
unsigned char *h_img, *h_back, *h_want;
unsigned char *d_img;
size_t pitch = 0;
int bad = 0, row, col;
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_img = (unsigned char *)malloc(w * h);
h_back = (unsigned char *)malloc(w * h);
h_want = (unsigned char *)malloc(w * h);
if (!h_img || !h_back || !h_want)
return 1;
CHECK(cudaMallocPitch(&d_img, &pitch, w, h));
printf("PITCH pitch=%zu\n", pitch);
printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img,
(void *)h_back, (void *)h_want, (void *)d_img);
/* P1: full cudaMemset over the pitched allocation, 2D readback. */
CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h));
memset(h_back, 0, w * h);
CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
cudaMemcpyDeviceToHost));
memset(h_want, FILL_BYTE, w * h);
bad += phase_check("P1", h_back, h_want, w * h, 1);
/* P2: pitched H2D pattern -> D2H round trip. */
for (row = 0; row < (int)h; row++)
for (col = 0; col < (int)w; col++)
h_img[row * w + col] = (unsigned char)((row + col) & 0xFF);
CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h,
cudaMemcpyHostToDevice));
memset(h_back, 0, w * h);
CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
cudaMemcpyDeviceToHost));
bad += phase_check("P2", h_back, h_img, w * h, 2);
/* P3: cudaMemset2D of the upper-left quadrant only. */
CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2));
for (row = 0; row < (int)h; row++)
for (col = 0; col < (int)w; col++) {
unsigned char expect = h_img[row * w + col];
if (row < (int)(h / 2) && col < (int)(w / 2))
expect = QUAD_BYTE;
h_want[row * w + col] = expect;
}
memset(h_back, 0, w * h);
CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
cudaMemcpyDeviceToHost));
bad += phase_check("P3", h_back, h_want, w * h, 3);
printf("MEMSET2D2 bad_total=%d\n", bad);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
printf("STATUS OK\n");
return 0;
}
|
