1 | // RUN: %libomptarget-compile-run-and-check-generic |
2 | // REQUIRES: ompt |
3 | // REQUIRES: gpu |
4 | |
5 | /* |
6 | * Example OpenMP program that registers EMI callbacks. |
7 | * Explicitly testing for an initialized device num and |
8 | * #pragma omp target [data enter / data exit / update] |
9 | * The latter with the addition of a nowait clause. |
10 | */ |
11 | |
12 | #include <omp.h> |
13 | #include <stdio.h> |
14 | |
15 | #include "callbacks.h" |
16 | #include "register_emi.h" |
17 | |
18 | #define N 100000 |
19 | |
20 | #pragma omp declare target |
21 | int c[N]; |
22 | #pragma omp end declare target |
23 | |
24 | int main() { |
25 | int a[N]; |
26 | int b[N]; |
27 | |
28 | int i; |
29 | |
30 | for (i = 0; i < N; i++) |
31 | a[i] = 0; |
32 | |
33 | for (i = 0; i < N; i++) |
34 | b[i] = i; |
35 | |
36 | for (i = 0; i < N; i++) |
37 | c[i] = 0; |
38 | |
39 | #pragma omp target enter data map(to : a) |
40 | #pragma omp target parallel for |
41 | { |
42 | for (int j = 0; j < N; j++) |
43 | a[j] = b[j]; |
44 | } |
45 | #pragma omp target exit data map(from : a) |
46 | |
47 | #pragma omp target parallel for map(alloc : c) |
48 | { |
49 | for (int j = 0; j < N; j++) |
50 | c[j] = 2 * j + 1; |
51 | } |
52 | #pragma omp target update from(c) nowait |
53 | #pragma omp barrier |
54 | |
55 | int rc = 0; |
56 | for (i = 0; i < N; i++) { |
57 | if (a[i] != i) { |
58 | rc++; |
59 | printf(format: "Wrong value: a[%d]=%d\n" , i, a[i]); |
60 | } |
61 | } |
62 | |
63 | for (i = 0; i < N; i++) { |
64 | if (c[i] != 2 * i + 1) { |
65 | rc++; |
66 | printf(format: "Wrong value: c[%d]=%d\n" , i, c[i]); |
67 | } |
68 | } |
69 | |
70 | if (!rc) |
71 | printf(format: "Success\n" ); |
72 | |
73 | return rc; |
74 | } |
75 | |
76 | /// CHECK-NOT: Callback Target EMI: |
77 | /// CHECK-NOT: device_num=-1 |
78 | /// CHECK: Callback Init: |
79 | /// CHECK: Callback Load: |
80 | /// CHECK: Callback Target EMI: kind=2 endpoint=1 |
81 | /// CHECK-NOT: device_num=-1 |
82 | /// CHECK-NOT: code=(nil) |
83 | /// CHECK: code=[[CODE1:.*]] |
84 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=1 |
85 | /// CHECK: code=[[CODE1]] |
86 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=1 |
87 | /// CHECK-NOT: dest=(nil) |
88 | /// CHECK: code=[[CODE1]] |
89 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=2 |
90 | /// CHECK: code=[[CODE1]] |
91 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=2 |
92 | /// CHECK: code=[[CODE1]] |
93 | /// CHECK: Callback Target EMI: kind=2 endpoint=2 |
94 | /// CHECK-NOT: device_num=-1 |
95 | /// CHECK: code=[[CODE1]] |
96 | /// CHECK: Callback Target EMI: kind=1 endpoint=1 |
97 | /// CHECK-NOT: device_num=-1 |
98 | /// CHECK-NOT: code=(nil) |
99 | /// CHECK: code=[[CODE2:.*]] |
100 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=1 |
101 | /// CHECK: code=[[CODE2]] |
102 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=1 |
103 | /// CHECK-NOT: dest=(nil) |
104 | /// CHECK: code=[[CODE2]] |
105 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=2 |
106 | /// CHECK: code=[[CODE2]] |
107 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=2 |
108 | /// CHECK: code=[[CODE2]] |
109 | /// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1 |
110 | /// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1 |
111 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=3 |
112 | /// CHECK: code=[[CODE2]] |
113 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=3 |
114 | /// CHECK: code=[[CODE2]] |
115 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=4 |
116 | /// CHECK: code=[[CODE2]] |
117 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=4 |
118 | /// CHECK: code=[[CODE2]] |
119 | /// CHECK: Callback Target EMI: kind=1 endpoint=2 |
120 | /// CHECK-NOT: device_num=-1 |
121 | /// CHECK: code=[[CODE2]] |
122 | /// CHECK: Callback Target EMI: kind=3 endpoint=1 |
123 | /// CHECK-NOT: device_num=-1 |
124 | /// CHECK-NOT: code=(nil) |
125 | /// CHECK: code=[[CODE3:.*]] |
126 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=3 |
127 | /// CHECK: code=[[CODE3]] |
128 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=3 |
129 | /// CHECK: code=[[CODE3]] |
130 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=4 |
131 | /// CHECK: code=[[CODE3]] |
132 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=4 |
133 | /// CHECK: code=[[CODE3]] |
134 | /// CHECK: Callback Target EMI: kind=3 endpoint=2 |
135 | /// CHECK-NOT: device_num=-1 |
136 | /// CHECK: code=[[CODE3]] |
137 | /// CHECK: Callback Target EMI: kind=1 endpoint=1 |
138 | /// CHECK-NOT: device_num=-1 |
139 | /// CHECK-NOT: code=(nil) |
140 | /// CHECK: code=[[CODE4:.*]] |
141 | /// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1 |
142 | /// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1 |
143 | /// CHECK: Callback Target EMI: kind=1 endpoint=2 |
144 | /// CHECK-NOT: device_num=-1 |
145 | /// CHECK: code=[[CODE4]] |
146 | /// CHECK: Callback Target EMI: kind=4 endpoint=1 |
147 | /// CHECK-NOT: device_num=-1 |
148 | /// CHECK-NOT: code=(nil) |
149 | /// CHECK: code=[[CODE5:.*]] |
150 | /// CHECK: Callback DataOp EMI: endpoint=1 optype=3 |
151 | /// CHECK: code=[[CODE5]] |
152 | /// CHECK: Callback DataOp EMI: endpoint=2 optype=3 |
153 | /// CHECK: code=[[CODE5]] |
154 | /// CHECK: Callback Target EMI: kind=4 endpoint=2 |
155 | /// CHECK-NOT: device_num=-1 |
156 | /// CHECK: code=[[CODE5]] |
157 | /// CHECK: Callback Fini: |
158 | |