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
main()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("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("Wrong value: c[%d]=%d\n", i, c[i]);
67 }
68 }
69
70 if (!rc)
71 printf("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