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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* p2_streams - three created streams, two kernels per stream (scale then
* add), independent 1024-float buffers, per-stream events.
*
* Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32).
* After all streams complete, each event must query as complete.
* Device memory: 3 x 4 KB (3 pages).
*/
#include "common.h"
#define SN 1024
#define NS 3
__global__ void scale(float *x, float s, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
x[i] = x[i] * s;
}
__global__ void addv(float *x, float v, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
x[i] = x[i] + v;
}
static void fill(float *h, int n)
{
for (int i = 0; i < n; i++)
h[i] = (float)(i % 100);
}
static float expected1(int i)
{
return ((float)(i % 100) * 2.0f) + 1.0f;
}
int main(int argc, char **argv)
{
float *h = (float *)malloc(sizeof(float) * SN);
float *got = (float *)malloc(sizeof(float) * SN);
int bad = 0;
if (!h || !got)
return 1;
fill(h, SN);
if (p2_emit_mode(argc, argv)) {
printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n",
NS, NS * 2, SN);
printf("RESULT streams events=%d bad=0\n", NS);
free(h); free(got);
return p2_status(0);
}
p2_device_info();
cudaStream_t st[NS];
cudaEvent_t ev[NS];
float *d_x[NS];
for (int k = 0; k < NS; k++) {
P2_CHECK(cudaStreamCreate(&st[k]));
P2_CHECK(cudaEventCreate(&ev[k]));
P2_CHECK(cudaMalloc(&d_x[k], sizeof(float) * SN));
P2_CHECK(cudaMemcpyAsync(d_x[k], h, sizeof(float) * SN,
cudaMemcpyHostToDevice, st[k]));
scale<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 2.0f, SN);
addv<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 1.0f, SN);
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaEventRecord(ev[k], st[k]));
}
int ibad = 0;
for (int k = 0; k < NS; k++) {
cudaError_t q;
int sbad = 0;
P2_CHECK(cudaStreamSynchronize(st[k]));
P2_CHECK(cudaMemcpy(got, d_x[k], sizeof(float) * SN,
cudaMemcpyDeviceToHost));
for (int i = 0; i < SN; i++)
if (got[i] != expected1(i))
sbad++;
q = cudaEventQuery(ev[k]);
if (q != cudaSuccess)
sbad++;
printf("RESULT streams stream=%d query_rc=%d bad=%d\n", k, (int)q,
sbad);
ibad += sbad;
}
printf("RESULT streams nstreams=%d launches=%d n=%d bad=%d\n",
NS, NS * 2, SN, ibad);
printf("RESULT streams events=%d bad=%d\n", NS, ibad == 0 ? 0 : 1);
bad = ibad;
for (int k = 0; k < NS; k++) {
cudaEventDestroy(ev[k]);
cudaStreamDestroy(st[k]);
cudaFree(d_x[k]);
}
free(h); free(got);
return p2_status(bad);
}
|
