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
13static void check(char *X, int Dev) {
14 printf(format: " host X = %c\n", *X);
15 char DV = -1;
16#pragma omp target device(Dev) map(from : DV)
17 DV = *X;
18 printf(format: "device X = %c\n", DV);
19}
20
21#define CHECK_DATA() check(&X, DevDefault)
22
23int 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(format: "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(format: "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(format: "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(format: "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

source code of offload/test/offloading/host_as_target.c