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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* ip_probe.cu - how many CTAs actually executed for an in-place kernel?
*
* scale_mark marks progress[blockIdx.x]=0xCAFEBABE after the c store; D2H of
* progress counts executed CTAs exactly. c (the in-place buffer) is read back
* and verified. seed (argv[1]) makes the run content-unique.
*
* Build: nvcc -O2 --cudart=static -o ip_probe ip_probe.cu
*/
#include <cstdio>
#include <cstdlib>
#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)
__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];
}
__global__ void scale_mark(float *c, float s, int n, unsigned *prog)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
c[i] = c[i] * s;
if (threadIdx.x == 0)
prog[blockIdx.x] = 0xCAFEBABEu;
}
int main(int argc, char **argv)
{
cudaDeviceProp prop;
const int n = 1 << 15;
const int grid = (n + 255) / 256;
size_t bytes = (size_t)n * sizeof(float);
float k = argc > 1 ? atof(argv[1]) : 0.0f;
float *ha, *hb, *hc;
float *da, *db, *dc;
unsigned *d_prog, h_prog[128];
int i;
setvbuf(stdout, NULL, _IONBF, 0);
CHECK(cudaGetDeviceProperties(&prop, 0));
printf("DEVICE %s seed=%g grid=%d\n", prop.name, (double)k, grid);
ha = (float *)malloc(bytes);
hb = (float *)malloc(bytes);
hc = (float *)malloc(bytes);
for (i = 0; i < n; i++) {
ha[i] = (float)(i % 1000) * 0.5f + k;
hb[i] = (float)(i % 997) * 0.25f + 2.0f * k;
}
memset(hc, 0xCD, bytes);
CHECK(cudaMalloc(&da, bytes));
CHECK(cudaMalloc(&db, bytes));
CHECK(cudaMalloc(&dc, bytes));
CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned)));
printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db,
(void *)dc, (void *)d_prog);
CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned)));
vadd<<<grid, 256>>>(da, db, dc, n);
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
{
int b = 0;
for (i = 0; i < n; i++)
if (hc[i] != ha[i] + hb[i])
b++;
printf("VADD readback bad=%d/%d\n", b, n);
}
/* in-place scale with CTA markers */
scale_mark<<<grid, 256>>>(dc, 2.0f, n, d_prog);
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned),
cudaMemcpyDeviceToHost));
{
int marked = 0, first_missing = -1;
for (i = 0; i < grid; i++) {
if (h_prog[i] == 0xCAFEBABEu)
marked++;
else if (first_missing < 0)
first_missing = i;
}
printf("CTA-MARK marked=%d/%d first_missing=%d\n", marked, grid,
first_missing);
}
CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
{
int b = 0, first = -1;
for (i = 0; i < n; i++)
if (hc[i] != (ha[i] + hb[i]) * 2.0f) {
b++;
if (first < 0)
first = i;
}
printf("SCALE readback bad=%d/%d first=%d\n", b, n, first);
if (first >= 0)
printf("SCALE firstbad got=%g want=%g\n", hc[first],
(ha[first] + hb[first]) * 2.0f);
}
printf("STATUS DONE\n");
return 0;
}
|
