1 // Check that specifying device as omp_get_initial_device():
2 // - Doesn't cause the runtime to fail.
3 // - Offloads code to the host.
4 // - Doesn't transfer data. In this case, just check that neither host data nor
5 // default device data are affected by the specified transfers.
6 // - Works whether it's specified directly or as the default device.
7
8 // RUN: %libomptarget-compile-run-and-check-generic
9
10 #include <omp.h>
11 #include <stdio.h>
12
check(char * X,int Dev)13 static void check(char *X, int Dev) {
14 printf(" host X = %c\n", *X);
15 char DV = -1;
16 #pragma omp target device(Dev) map(from : DV)
17 DV = *X;
18 printf("device X = %c\n", DV);
19 }
20
21 #define CHECK_DATA() check(&X, DevDefault)
22
main(void)23 int main(void) {
24 int DevDefault = omp_get_default_device();
25 int DevInit = omp_get_initial_device();
26
27 //--------------------------------------------------
28 // Initialize data on the host and default device.
29 //--------------------------------------------------
30
31 // CHECK: host X = h
32 // CHECK-NEXT: device X = d
33 char X = 'd';
34 #pragma omp target enter data map(to : X)
35 X = 'h';
36 CHECK_DATA();
37
38 //--------------------------------------------------
39 // Check behavior when specifying host directly.
40 //--------------------------------------------------
41
42 // CHECK-NEXT: omp_is_initial_device() = 1
43 // CHECK-NEXT: host X = h
44 // CHECK-NEXT: device X = d
45 #pragma omp target device(DevInit) map(always, tofrom : X)
46 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
47 CHECK_DATA();
48
49 // CHECK-NEXT: omp_is_initial_device() = 1
50 // CHECK-NEXT: host X = h
51 // CHECK-NEXT: device X = d
52 #pragma omp target teams device(DevInit) num_teams(1) map(always, tofrom : X)
53 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
54 CHECK_DATA();
55
56 // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
57 // how to check that it actually pushes to the initial device.
58 #pragma omp target teams device(DevInit) num_teams(1)
59 #pragma omp distribute
60 for (int i = 0; i < 2; ++i)
61 ;
62
63 // CHECK-NEXT: host X = h
64 // CHECK-NEXT: device X = d
65 #pragma omp target data device(DevInit) map(always, tofrom : X)
66 ;
67 CHECK_DATA();
68
69 // CHECK-NEXT: host X = h
70 // CHECK-NEXT: device X = d
71 #pragma omp target enter data device(DevInit) map(always, to : X)
72 ;
73 CHECK_DATA();
74
75 // CHECK-NEXT: host X = h
76 // CHECK-NEXT: device X = d
77 #pragma omp target exit data device(DevInit) map(always, from : X)
78 ;
79 CHECK_DATA();
80
81 // CHECK-NEXT: host X = h
82 // CHECK-NEXT: device X = d
83 #pragma omp target update device(DevInit) to(X)
84 ;
85 CHECK_DATA();
86
87 // CHECK-NEXT: host X = h
88 // CHECK-NEXT: device X = d
89 #pragma omp target update device(DevInit) from(X)
90 ;
91 CHECK_DATA();
92
93 //--------------------------------------------------
94 // Check behavior when device defaults to host.
95 //--------------------------------------------------
96
97 omp_set_default_device(DevInit);
98
99 // CHECK-NEXT: omp_is_initial_device() = 1
100 // CHECK-NEXT: host X = h
101 // CHECK-NEXT: device X = d
102 #pragma omp target map(always, tofrom : X)
103 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
104 CHECK_DATA();
105
106 // CHECK-NEXT: omp_is_initial_device() = 1
107 // CHECK-NEXT: host X = h
108 // CHECK-NEXT: device X = d
109 #pragma omp target teams num_teams(1) map(always, tofrom : X)
110 printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
111 CHECK_DATA();
112
113 // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
114 // how to check that it actually pushes to the initial device.
115 #pragma omp target teams num_teams(1)
116 #pragma omp distribute
117 for (int i = 0; i < 2; ++i)
118 ;
119
120 // CHECK-NEXT: host X = h
121 // CHECK-NEXT: device X = d
122 #pragma omp target data map(always, tofrom : X)
123 ;
124 CHECK_DATA();
125
126 // CHECK-NEXT: host X = h
127 // CHECK-NEXT: device X = d
128 #pragma omp target enter data map(always, to : X)
129 ;
130 CHECK_DATA();
131
132 // CHECK-NEXT: host X = h
133 // CHECK-NEXT: device X = d
134 #pragma omp target exit data map(always, from : X)
135 ;
136 CHECK_DATA();
137
138 // CHECK-NEXT: host X = h
139 // CHECK-NEXT: device X = d
140 #pragma omp target update to(X)
141 ;
142 CHECK_DATA();
143
144 // CHECK-NEXT: host X = h
145 // CHECK-NEXT: device X = d
146 #pragma omp target update from(X)
147 ;
148 CHECK_DATA();
149
150 return 0;
151 }
152