1 /* Test dynamic refcount and copy behavior of separate structure members. */
2
3 #include <assert.h>
4 #include <stdbool.h>
5 #include <openacc.h>
6
7 struct s
8 {
9 signed char a;
10 float b;
11 };
12
13 static void test(unsigned variant)
14 {
15 struct s s = { .a = 73, .b = -22 };
16
17 #pragma acc enter data copyin(s.a, s.b)
18 assert(acc_is_present(&s.a, sizeof s.a));
19 assert(acc_is_present(&s.b, sizeof s.b));
20
21 /* To verify that any following 'copyin' doesn't 'copyin' again. */
22 s.a = -s.a;
23 s.b = -s.b;
24
25 if (variant & 4)
26 {
27 if (variant & 8)
28 {
29 #pragma acc enter data copyin(s.b)
30 }
31 else
32 acc_copyin(&s.b, sizeof s.b);
33 assert(acc_is_present(&s.a, sizeof s.a));
34 assert(acc_is_present(&s.b, sizeof s.b));
35
36 if (variant & 16)
37 {
38 #pragma acc enter data copyin(s.a)
39 }
40 else
41 acc_copyin(&s.a, sizeof s.a);
42 assert(acc_is_present(&s.a, sizeof s.a));
43 assert(acc_is_present(&s.b, sizeof s.b));
44
45 if (variant & 32)
46 {
47 #pragma acc enter data copyin(s.a)
48 acc_copyin(&s.b, sizeof s.b);
49 #pragma acc enter data copyin(s.b)
50 #pragma acc enter data copyin(s.b)
51 acc_copyin(&s.a, sizeof s.a);
52 acc_copyin(&s.a, sizeof s.a);
53 acc_copyin(&s.a, sizeof s.a);
54 }
55 assert(acc_is_present(&s.a, sizeof s.a));
56 assert(acc_is_present(&s.b, sizeof s.b));
57 }
58
59 #pragma acc parallel \
60 copy(s.a, s.b)
61 {
62 #if ACC_MEM_SHARED
63 if (s.a++ != -73)
64 __builtin_abort();
65 if (s.b-- != 22)
66 __builtin_abort();
67 #else
68 if (s.a++ != 73)
69 __builtin_abort();
70 if (s.b-- != -22)
71 __builtin_abort();
72 #endif
73 }
74 #if ACC_MEM_SHARED
75 assert(s.a == -72);
76 assert(s.b == 21);
77 #else
78 assert(s.a == -73);
79 assert(s.b == 22);
80 #endif
81
82 if (variant & 32)
83 {
84 if (variant & 1)
85 {
86 #pragma acc exit data copyout(s.a) finalize
87 }
88 else
89 acc_copyout_finalize(&s.a, sizeof s.a);
90 }
91 else
92 {
93 if (variant & 1)
94 {
95 #pragma acc exit data copyout(s.a)
96 }
97 else
98 acc_copyout(&s.a, sizeof s.a);
99 if (variant & 4)
100 {
101 assert(acc_is_present(&s.a, sizeof s.a));
102 assert(acc_is_present(&s.b, sizeof s.b));
103 #if ACC_MEM_SHARED
104 assert(s.a == -72);
105 assert(s.b == 21);
106 #else
107 assert(s.a == -73);
108 assert(s.b == 22);
109 #endif
110 if (variant & 1)
111 {
112 #pragma acc exit data copyout(s.a)
113 }
114 else
115 acc_copyout(&s.a, sizeof s.a);
116 }
117 }
118 #if ACC_MEM_SHARED
119 assert(acc_is_present(&s.a, sizeof s.a));
120 assert(acc_is_present(&s.b, sizeof s.b));
121 assert(s.a == -72);
122 assert(s.b == 21);
123 #else
124 assert(!acc_is_present(&s.a, sizeof s.a));
125 assert(acc_is_present(&s.b, sizeof s.b));
126 assert(s.a == 74);
127 assert(s.b == 22);
128 #endif
129
130 if (variant & 32)
131 {
132 if (variant & 2)
133 {
134 #pragma acc exit data copyout(s.b) finalize
135 }
136 else
137 acc_copyout_finalize(&s.b, sizeof s.b);
138 }
139 else
140 {
141 if (variant & 2)
142 {
143 #pragma acc exit data copyout(s.b)
144 }
145 else
146 acc_copyout(&s.b, sizeof s.b);
147 if (variant & 4)
148 {
149 #if ACC_MEM_SHARED
150 assert(acc_is_present(&s.a, sizeof s.a));
151 assert(acc_is_present(&s.b, sizeof s.b));
152 assert(s.a == -72);
153 assert(s.b == 21);
154 #else
155 assert(!acc_is_present(&s.a, sizeof s.a));
156 assert(acc_is_present(&s.b, sizeof s.b));
157 assert(s.a == 74);
158 assert(s.b == 22);
159 #endif
160 if (variant & 2)
161 {
162 #pragma acc exit data copyout(s.b)
163 }
164 else
165 acc_copyout(&s.b, sizeof s.b);
166 }
167 }
168 #if ACC_MEM_SHARED
169 assert(acc_is_present(&s.a, sizeof s.a));
170 assert(acc_is_present(&s.b, sizeof s.b));
171 assert(s.a == -72);
172 assert(s.b == 21);
173 #else
174 assert(!acc_is_present(&s.a, sizeof s.a));
175 assert(!acc_is_present(&s.b, sizeof s.b));
176 assert(s.a == 74);
177 assert(s.b == -23);
178 #endif
179 }
180
181 int main()
182 {
183 for (unsigned variant = 0; variant < 64; ++variant)
184 test(variant);
185
186 return 0;
187 }