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