Skip to content

Commit 57f0bdc

Browse files
mhalkmhalk
committed
[OpenMP][OMPT] Fix target enter data callback ordering & reported device num
This patch fixes: #64738 We observed multiple issues, primarily that the `DeviceId` was reported as -1 in certain scenarios. The reason for this is simply that the device is not initialized at that point. Hence, we need to move the RAII object creation just after the `checkDeviceAndCtors`, closer to the actual call we want to observe. This also solves an odering issue where one `target enter data` callback would be executed before the `Init` callback. Additionally, this change will also fix that the callbacks corresponding to `enter / exit data` and `update` in conjunction with `nowait` would not result in the emission of an OMPT callback. Added a testcase to cover initialized device number and `omp target` constructs. Reviewed By: dhruvachak Differential Revision: https://reviews.llvm.org/D157605
1 parent a886870 commit 57f0bdc

File tree

2 files changed

+147
-20
lines changed

2 files changed

+147
-20
lines changed

Diff for: ‎openmp/libomptarget/src/interface.cpp

+19-20
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,21 @@ targetDataMapper(ident_t *Loc, int64_t DeviceId, int32_t ArgNum,
108108
TargetAsyncInfoTy TargetAsyncInfo(Device);
109109
AsyncInfoTy &AsyncInfo = TargetAsyncInfo;
110110

111+
/// RAII to establish tool anchors before and after data begin / end / update
112+
OMPT_IF_BUILT(assert((TargetDataFunction == targetDataBegin ||
113+
TargetDataFunction == targetDataEnd ||
114+
TargetDataFunction == targetDataUpdate) &&
115+
"Encountered unexpected TargetDataFunction during "
116+
"execution of targetDataMapper");
117+
auto CallbackFunctions =
118+
(TargetDataFunction == targetDataBegin)
119+
? RegionInterface.getCallbacks<ompt_target_enter_data>()
120+
: (TargetDataFunction == targetDataEnd)
121+
? RegionInterface.getCallbacks<ompt_target_exit_data>()
122+
: RegionInterface.getCallbacks<ompt_target_update>();
123+
InterfaceRAII TargetDataRAII(CallbackFunctions, DeviceId,
124+
OMPT_GET_RETURN_ADDRESS(0));)
125+
111126
int Rc = OFFLOAD_SUCCESS;
112127
Rc = TargetDataFunction(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes,
113128
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
@@ -129,12 +144,6 @@ EXTERN void __tgt_target_data_begin_mapper(ident_t *Loc, int64_t DeviceId,
129144
map_var_info_t *ArgNames,
130145
void **ArgMappers) {
131146
TIMESCOPE_WITH_IDENT(Loc);
132-
/// RAII to establish tool anchors before and after data begin
133-
OMPT_IF_BUILT(InterfaceRAII TargetDataEnterRAII(
134-
RegionInterface.getCallbacks<ompt_target_enter_data>(),
135-
DeviceId,
136-
/* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
137-
138147
targetDataMapper<AsyncInfoTy>(Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes,
139148
ArgTypes, ArgNames, ArgMappers, targetDataBegin,
140149
"Entering OpenMP data region", "begin");
@@ -161,12 +170,6 @@ EXTERN void __tgt_target_data_end_mapper(ident_t *Loc, int64_t DeviceId,
161170
map_var_info_t *ArgNames,
162171
void **ArgMappers) {
163172
TIMESCOPE_WITH_IDENT(Loc);
164-
/// RAII to establish tool anchors before and after data end
165-
OMPT_IF_BUILT(InterfaceRAII TargetDataExitRAII(
166-
RegionInterface.getCallbacks<ompt_target_exit_data>(),
167-
DeviceId,
168-
/* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
169-
170173
targetDataMapper<AsyncInfoTy>(Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes,
171174
ArgTypes, ArgNames, ArgMappers, targetDataEnd,
172175
"Exiting OpenMP data region", "end");
@@ -190,12 +193,6 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *Loc, int64_t DeviceId,
190193
map_var_info_t *ArgNames,
191194
void **ArgMappers) {
192195
TIMESCOPE_WITH_IDENT(Loc);
193-
/// RAII to establish tool anchors before and after data update
194-
OMPT_IF_BUILT(InterfaceRAII TargetDataUpdateRAII(
195-
RegionInterface.getCallbacks<ompt_target_update>(),
196-
DeviceId,
197-
/* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
198-
199196
targetDataMapper<AsyncInfoTy>(
200197
Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames,
201198
ArgMappers, targetDataUpdate, "Updating OpenMP data", "update");
@@ -295,7 +292,8 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
295292
DeviceTy &Device = *PM->Devices[DeviceId];
296293
TargetAsyncInfoTy TargetAsyncInfo(Device);
297294
AsyncInfoTy &AsyncInfo = TargetAsyncInfo;
298-
OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII(
295+
/// RAII to establish tool anchors before and after target region
296+
OMPT_IF_BUILT(InterfaceRAII TargetRAII(
299297
RegionInterface.getCallbacks<ompt_target>(), DeviceId,
300298
/* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
301299

@@ -386,7 +384,8 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
386384
return OMP_TGT_FAIL;
387385
}
388386
DeviceTy &Device = *PM->Devices[DeviceId];
389-
OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII(
387+
/// RAII to establish tool anchors before and after target region
388+
OMPT_IF_BUILT(InterfaceRAII TargetRAII(
390389
RegionInterface.getCallbacks<ompt_target>(), DeviceId,
391390
/* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));)
392391

Diff for: ‎openmp/libomptarget/test/ompt/veccopy_data.c

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

0 commit comments

Comments
 (0)