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