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