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