]>
Commit | Line | Data |
---|---|---|
2620c80d TS |
1 | /* Test valid use of host_data directive. */ |
2 | ||
b3b75e66 TS |
3 | /* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ |
4 | ||
2620c80d TS |
5 | int v1[3][3]; |
6 | ||
7 | void | |
8 | f (void) | |
9 | { | |
10 | #pragma acc host_data use_device(v1) | |
b3b75e66 TS |
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" } } */ | |
2620c80d | 13 | ; |
d5c23c6c TB |
14 | |
15 | #pragma acc host_data use_device(v1) if_present | |
b3b75e66 TS |
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" } } */ | |
d5c23c6c | 18 | ; |
2620c80d TS |
19 | } |
20 | ||
d902b330 JJ |
21 | |
22 | void bar (float *, float *); | |
23 | ||
24 | void | |
b3b75e66 | 25 | foo (float *x, float *y, float *yy) |
d902b330 JJ |
26 | { |
27 | int n = 1 << 10; | |
d5c23c6c | 28 | #pragma acc data create(x[0:n]) |
d902b330 | 29 | { |
d5c23c6c TB |
30 | bar (x, y); |
31 | ||
32 | /* This should fail at run time because y is not mapped. */ | |
d902b330 | 33 | #pragma acc host_data use_device(x,y) |
b3b75e66 TS |
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" } } */ | |
d902b330 | 36 | bar (x, y); |
d5c23c6c TB |
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 | |
b3b75e66 TS |
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" } } */ | |
d5c23c6c TB |
43 | bar (x, y); |
44 | ||
b3b75e66 | 45 | #pragma acc data copyout(yy[0:n]) |
d5c23c6c | 46 | { |
b3b75e66 TS |
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); | |
d5c23c6c | 51 | |
b3b75e66 TS |
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); | |
d5c23c6c | 56 | |
b3b75e66 TS |
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); | |
d5c23c6c | 61 | |
b3b75e66 TS |
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); | |
d5c23c6c | 66 | } |
d902b330 JJ |
67 | } |
68 | } |