1 /* Test valid use of host_data directive. */
2
3 /* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
4
5 int v1[3][3];
6
7 void
f(void)8 f (void)
9 {
10 #pragma acc host_data use_device(v1)
11 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(v1\\)$" 1 "original" } }
12 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(v1\\)$" 1 "gimple" } } */
13 ;
14
15 #pragma acc host_data use_device(v1) if_present
16 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(v1\\)$" 1 "original" } }
17 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:v1\\)$" 1 "gimple" } } */
18 ;
19 }
20
21
22 void bar (float *, float *);
23
24 void
foo(float * x,float * y,float * yy)25 foo (float *x, float *y, float *yy)
26 {
27 int n = 1 << 10;
28 #pragma acc data create(x[0:n])
29 {
30 bar (x, y);
31
32 /* This should fail at run time because y is not mapped. */
33 #pragma acc host_data use_device(x,y)
34 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
35 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
36 bar (x, y);
37
38 /* y is still not mapped, but this should not fail at run time but
39 continue execution with y remaining as the host address. */
40 #pragma acc host_data use_device(x,y) if_present
41 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
42 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:y\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
43 bar (x, y);
44
45 #pragma acc data copyout(yy[0:n])
46 {
47 #pragma acc host_data use_device(x,yy)
48 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
49 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
50 bar (x, yy);
51
52 #pragma acc host_data use_device(x,yy) if_present
53 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
54 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
55 bar (x, yy);
56
57 #pragma acc host_data use_device(x,yy) if(x != yy)
58 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x \\!= yy\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
59 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
60 bar (x, yy);
61
62 #pragma acc host_data use_device(x,yy) if_present if(x == yy)
63 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x == yy\\) if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
64 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
65 bar (x, yy);
66 }
67 }
68 }
69