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
|
/* Verify back to back 'async' operations, two data mappings.
Make sure that despite two data mappings, this isn't using the libgomp
'cbuf' buffering.
*/
#include <stdlib.h>
#define N 128
static void
t1 (void)
{
unsigned int *a, *b;
int i;
int nbytes;
nbytes = N * sizeof (unsigned int);
a = (unsigned int *) malloc (nbytes);
b = (unsigned int *) malloc (nbytes);
for (i = 0; i < N; i++)
b[i] = a[i] = 3;
#pragma acc parallel async copy (a[0:N], b[0:N])
for (int ii = 0; ii < N; ii++)
b[ii] += (a[ii] += 1);
#pragma acc parallel async copy (a[0:N], b[0:N])
for (int ii = 0; ii < N; ii++)
b[ii] += (a[ii] += 1);
#pragma acc wait
for (i = 0; i < N; i++)
{
if (a[i] != 5)
abort ();
if (b[i] != 12)
abort ();
}
}
static void
t2 (void)
{
unsigned int *a, *b;
int i;
int nbytes;
nbytes = N * sizeof (unsigned int);
a = (unsigned int *) malloc (nbytes);
b = (unsigned int *) malloc (nbytes);
#pragma acc data copyin (a[0:N], b[0:N])
{
for (i = 0; i < N; i++)
b[i] = a[i] = 3;
#pragma acc update async device (a[0:N], b[0:N])
#pragma acc parallel async present (a[0:N], b[0:N])
for (int ii = 0; ii < N; ii++)
b[ii] += (a[ii] += 1);
#pragma acc update async host (a[0:N], b[0:N])
#pragma acc update async device (a[0:N], b[0:N])
#pragma acc parallel async present (a[0:N], b[0:N])
for (int ii = 0; ii < N; ii++)
b[ii] += (a[ii] += 1);
#pragma acc update async host (a[0:N], b[0:N])
#pragma acc wait
}
for (i = 0; i < N; i++)
{
if (a[i] != 5)
abort ();
if (b[i] != 12)
abort ();
}
}
int
main (void)
{
t1 ();
t2 ();
return 0;
}
|
