Skip to content

Commit 48421ac

Browse files
jdenny-ornljdenny-ornl
committed
[OpenMP] Improve ref count debug messages
For example, without this patch: ``` $ cat test.c int main() { int x; #pragma omp target enter data map(alloc: x) #pragma omp target exit data map(release: x) ; return 0; } $ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda test.c $ LIBOMPTARGET_DEBUG=1 ./a.out |& grep 'Creating\|Mapping exists' Libomptarget --> Creating new map entry with HstPtrBegin=0x00007ffcace8e448, TgtPtrBegin=0x00007f12ef600000, Size=4, Name=unknown Libomptarget --> Mapping exists with HstPtrBegin=0x00007ffcace8e448, TgtPtrBegin=0x00007f12ef600000, Size=4, updated RefCount=1 ``` There are two problems in this example: * `RefCount` is not reported when a mapping is created, but it might be 1 or infinite. In this case, because it's created by `omp target enter data`, it's 1. Seeing that would make later `RefCount` messages easier to understand. * `RefCount` is still 1 at the `omp target exit data`, but it's reported as `updated`. The reason it's still 1 is that, upon deletions, the reference count is generally not updated in `DeviceTy::getTgtPtrBegin`, where the report is produced. Instead, it's zeroed later in `DeviceTy::deallocTgtPtr`, where it's actually removed from the mapping table. This patch makes the following changes: * Report the reference count when creating a mapping. * Where an existing mapping is reported, always report a reference count action: * `update suppressed` when `UpdateRefCount=false` * `incremented` * `decremented` * `deferred final decrement`, which replaces the misleading `updated` in the above example * Add comments to `DeviceTy::getTgtPtrBegin` to explain why it does not zero the reference count. (Please advise if these comments miss the point.) * For unified shared memory, don't report confusing messages like `RefCount=` or `RefCount= updated` given that reference counts are irrelevant in this case. Instead, just report `for unified shared memory`. * Use `INFO` not `DP` consistently for `Mapping exists` messages. * Fix device table dumps to print `INF` instead of `-1` for an infinite reference count. Reviewed By: jhuber6, grokos Differential Revision: https://reviews.llvm.org/D104559
1 parent 0c0628c commit 48421ac

File tree

5 files changed

+112
-68
lines changed

5 files changed

+112
-68
lines changed

‎openmp/docs/design/Runtimes.rst‎

Lines changed: 47 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -141,50 +141,66 @@ provide the following output from the runtime library.
141141
Info: Entering OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
142142
Info: to(X[0:N])[16384]
143143
Info: tofrom(Y[0:N])[16384]
144-
Info: Creating new map entry with HstPtrBegin=0x00007ffde9e99000,
145-
TgtPtrBegin=0x00007f15dc600000, Size=16384, Name=X[0:N]
146-
Info: Copying data from host to device, HstPtr=0x00007ffde9e99000,
147-
TgtPtr=0x00007f15dc600000, Size=16384, Name=X[0:N]
148-
Info: Creating new map entry with HstPtrBegin=0x00007ffde9e95000,
149-
TgtPtrBegin=0x00007f15dc604000, Size=16384, Name=Y[0:N]
150-
Info: Copying data from host to device, HstPtr=0x00007ffde9e95000,
151-
TgtPtr=0x00007f15dc604000, Size=16384, Name=Y[0:N]
144+
Info: Creating new map entry with HstPtrBegin=0x00007fff0d259a40,
145+
TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1, Name=X[0:N]
146+
Info: Copying data from host to device, HstPtr=0x00007fff0d259a40,
147+
TgtPtr=0x00007fdba5800000, Size=16384, Name=X[0:N]
148+
Info: Creating new map entry with HstPtrBegin=0x00007fff0d255a40,
149+
TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1, Name=Y[0:N]
150+
Info: Copying data from host to device, HstPtr=0x00007fff0d255a40,
151+
TgtPtr=0x00007fdba5804000, Size=16384, Name=Y[0:N]
152152
Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:14:1:
153153
Info: Host Ptr Target Ptr Size (B) RefCount Declaration
154-
Info: 0x00007ffde9e95000 0x00007f15dc604000 16384 1 Y[0:N] at zaxpy.cpp:13:17
155-
Info: 0x00007ffde9e99000 0x00007f15dc600000 16384 1 X[0:N] at zaxpy.cpp:13:11
154+
Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
155+
Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
156156
Info: Entering OpenMP kernel at zaxpy.cpp:6:1 with 4 arguments:
157157
Info: firstprivate(N)[8] (implicit)
158158
Info: use_address(Y)[0] (implicit)
159159
Info: tofrom(D)[16] (implicit)
160160
Info: use_address(X)[0] (implicit)
161-
Info: Mapping exists (implicit) with HstPtrBegin=0x00007ffde9e95000,
162-
TgtPtrBegin=0x00007f15dc604000, Size=0, updated RefCount=2, Name=Y
163-
Info: Creating new map entry with HstPtrBegin=0x00007ffde9e94fb0,
164-
TgtPtrBegin=0x00007f15dc608000, Size=16, Name=D
165-
Info: Copying data from host to device, HstPtr=0x00007ffde9e94fb0,
166-
TgtPtr=0x00007f15dc608000, Size=16, Name=D
167-
Info: Mapping exists (implicit) with HstPtrBegin=0x00007ffde9e99000,
168-
TgtPtrBegin=0x00007f15dc600000, Size=0, updated RefCount=2, Name=X
169-
Info: Launching kernel __omp_offloading_fd02_e25f6e76__Z5zaxpyPSt7complexIdES1_S0_m_l6
161+
Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d255a40,
162+
TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (incremented), Name=Y
163+
Info: Creating new map entry with HstPtrBegin=0x00007fff0d2559f0,
164+
TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1, Name=D
165+
Info: Copying data from host to device, HstPtr=0x00007fff0d2559f0,
166+
TgtPtr=0x00007fdba5808000, Size=16, Name=D
167+
Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d259a40,
168+
TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (incremented), Name=X
169+
Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
170+
TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (update suppressed)
171+
Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
172+
TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (update suppressed)
173+
Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
174+
TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (update suppressed)
175+
Info: Launching kernel __omp_offloading_10305_c08c86__Z5zaxpyPSt7complexIdES1_S0_m_l6
170176
with 8 blocks and 128 threads in SPMD mode
171-
Info: Copying data from device to host, TgtPtr=0x00007f15dc608000,
172-
HstPtr=0x00007ffde9e94fb0, Size=16, Name=D
173-
Info: Removing map entry with HstPtrBegin=0x00007ffde9e94fb0,
174-
TgtPtrBegin=0x00007f15dc608000, Size=16, Name=D
177+
Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
178+
TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=1 (decremented)
179+
Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
180+
TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (deferred final decrement)
181+
Info: Copying data from device to host, TgtPtr=0x00007fdba5808000,
182+
HstPtr=0x00007fff0d2559f0, Size=16, Name=D
183+
Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
184+
TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=1 (decremented)
185+
Info: Removing map entry with HstPtrBegin=0x00007fff0d2559f0,
186+
TgtPtrBegin=0x00007fdba5808000, Size=16, Name=D
175187
Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:6:1:
176188
Info: Host Ptr Target Ptr Size (B) RefCount Declaration
177-
Info: 0x00007ffde9e95000 0x00007f15dc604000 16384 1 Y[0:N] at zaxpy.cpp:13:17
178-
Info: 0x00007ffde9e99000 0x00007f15dc600000 16384 1 X[0:N] at zaxpy.cpp:13:11
189+
Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
190+
Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
179191
Info: Exiting OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
180192
Info: to(X[0:N])[16384]
181193
Info: tofrom(Y[0:N])[16384]
182-
Info: Copying data from device to host, TgtPtr=0x00007f15dc604000,
183-
HstPtr=0x00007ffde9e95000, Size=16384, Name=Y[0:N]
184-
Info: Removing map entry with HstPtrBegin=0x00007ffde9e95000,
185-
TgtPtrBegin=0x00007f15dc604000, Size=16384, Name=Y[0:N]
186-
Info: Removing map entry with HstPtrBegin=0x00007ffde9e99000,
187-
TgtPtrBegin=0x00007f15dc600000, Size=16384, Name=X[0:N]
194+
Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
195+
TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1 (deferred final decrement)
196+
Info: Copying data from device to host, TgtPtr=0x00007fdba5804000,
197+
HstPtr=0x00007fff0d255a40, Size=16384, Name=Y[0:N]
198+
Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
199+
TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1 (deferred final decrement)
200+
Info: Removing map entry with HstPtrBegin=0x00007fff0d255a40,
201+
TgtPtrBegin=0x00007fdba5804000, Size=16384, Name=Y[0:N]
202+
Info: Removing map entry with HstPtrBegin=0x00007fff0d259a40,
203+
TgtPtrBegin=0x00007fdba5800000, Size=16384, Name=X[0:N]
188204
189205
From this information, we can see the OpenMP kernel being launched on the CUDA
190206
device with enough threads and blocks for all ``1024`` iterations of the loop in

‎openmp/libomptarget/src/device.cpp‎

Lines changed: 48 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -76,16 +76,20 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
7676
}
7777

7878
// Mapping does not exist, allocate it with refCount=INF
79-
auto Res = HostDataToTargetMap.emplace(
80-
(uintptr_t)HstPtrBegin /*HstPtrBase*/,
81-
(uintptr_t)HstPtrBegin /*HstPtrBegin*/,
82-
(uintptr_t)HstPtrBegin + Size /*HstPtrEnd*/,
83-
(uintptr_t)TgtPtrBegin /*TgtPtrBegin*/, nullptr, true /*IsRefCountINF*/);
84-
auto NewEntry = Res.first;
79+
const HostDataToTargetTy &newEntry =
80+
*HostDataToTargetMap
81+
.emplace(
82+
/*HstPtrBase=*/(uintptr_t)HstPtrBegin,
83+
/*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
84+
/*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
85+
/*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin, /*Name=*/nullptr,
86+
/*IsRefCountINF=*/true)
87+
.first;
8588
DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD
86-
", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n",
87-
DPxPTR(NewEntry->HstPtrBase), DPxPTR(NewEntry->HstPtrBegin),
88-
DPxPTR(NewEntry->HstPtrEnd), DPxPTR(NewEntry->TgtPtrBegin));
89+
", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", RefCount=%s\n",
90+
DPxPTR(newEntry.HstPtrBase), DPxPTR(newEntry.HstPtrBegin),
91+
DPxPTR(newEntry.HstPtrEnd), DPxPTR(newEntry.TgtPtrBegin),
92+
newEntry.refCountToStr().c_str());
8993

9094
DataMapMtx.unlock();
9195

@@ -211,18 +215,16 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
211215
((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && IsImplicit)) {
212216
auto &HT = *lr.Entry;
213217
IsNew = false;
214-
215218
if (UpdateRefCount)
216219
HT.incRefCount();
217-
218220
uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
219221
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
220222
"Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
221223
", "
222-
"Size=%" PRId64 ",%s RefCount=%s, Name=%s\n",
224+
"Size=%" PRId64 ", RefCount=%s (%s), Name=%s\n",
223225
(IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(tp),
224-
Size, (UpdateRefCount ? " updated" : ""),
225-
HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str(),
226+
Size, HT.refCountToStr().c_str(),
227+
UpdateRefCount ? "incremented" : "update suppressed",
226228
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
227229
rc = (void *)tp;
228230
} else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
@@ -246,9 +248,9 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
246248
// In addition to the mapping rules above, the close map modifier forces the
247249
// mapping of the variable to the device.
248250
if (Size) {
249-
DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " RefCount=%s\n",
250-
DPxPTR((uintptr_t)HstPtrBegin), Size,
251-
(UpdateRefCount ? " updated" : ""));
251+
DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
252+
"memory\n",
253+
DPxPTR((uintptr_t)HstPtrBegin), Size);
252254
IsHostPtr = true;
253255
rc = HstPtrBegin;
254256
}
@@ -263,13 +265,18 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
263265
// If it is not contained and Size > 0, we should create a new entry for it.
264266
IsNew = true;
265267
uintptr_t tp = (uintptr_t)allocData(Size, HstPtrBegin);
268+
const HostDataToTargetTy &newEntry =
269+
*HostDataToTargetMap
270+
.emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
271+
(uintptr_t)HstPtrBegin + Size, tp, HstPtrName)
272+
.first;
266273
INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
267274
"Creating new map entry with "
268-
"HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, Name=%s\n",
275+
"HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
276+
"RefCount=%s, Name=%s\n",
269277
DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
278+
newEntry.refCountToStr().c_str(),
270279
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
271-
HostDataToTargetMap.emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
272-
(uintptr_t)HstPtrBegin + Size, tp, HstPtrName);
273280
rc = (void *)tp;
274281
}
275282

@@ -292,25 +299,35 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
292299
if (lr.Flags.IsContained ||
293300
(!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
294301
auto &HT = *lr.Entry;
302+
// We do not decrement the reference count to zero here. deallocTgtPtr does
303+
// that atomically with removing the mapping. Otherwise, before this thread
304+
// removed the mapping in deallocTgtPtr, another thread could retrieve the
305+
// mapping, increment and decrement back to zero, and then both threads
306+
// would try to remove the mapping, resulting in a double free.
295307
IsLast = HT.getRefCount() == 1;
296-
297-
if (!IsLast && UpdateRefCount)
308+
const char *RefCountAction;
309+
if (!UpdateRefCount)
310+
RefCountAction = "update suppressed";
311+
else if (IsLast)
312+
RefCountAction = "deferred final decrement";
313+
else {
314+
RefCountAction = "decremented";
298315
HT.decRefCount();
299-
316+
}
300317
uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
301-
DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
302-
"Size=%" PRId64 ",%s RefCount=%s\n",
303-
DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
304-
(UpdateRefCount ? " updated" : ""),
305-
HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
318+
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
319+
"Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
320+
"Size=%" PRId64 ", RefCount=%s (%s)\n",
321+
DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.refCountToStr().c_str(),
322+
RefCountAction);
306323
rc = (void *)tp;
307324
} else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
308325
// If the value isn't found in the mapping and unified shared memory
309326
// is on then it means we have stumbled upon a value which we need to
310327
// use directly from the host.
311-
DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " RefCount=%s\n",
312-
DPxPTR((uintptr_t)HstPtrBegin), Size,
313-
(UpdateRefCount ? " updated" : ""));
328+
DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
329+
"memory\n",
330+
DPxPTR((uintptr_t)HstPtrBegin), Size);
314331
IsHostPtr = true;
315332
rc = HstPtrBegin;
316333
}

‎openmp/libomptarget/src/device.h‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,6 +88,10 @@ struct HostDataToTargetTy {
8888
}
8989

9090
bool isRefCountInf() const { return RefCount == INFRefCount; }
91+
92+
std::string refCountToStr() const {
93+
return isRefCountInf() ? "INF" : std::to_string(getRefCount());
94+
}
9195
};
9296

9397
typedef uintptr_t HstPtrBeginTy;

‎openmp/libomptarget/src/private.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -118,11 +118,11 @@ static inline void dumpTargetPointerMappings(const ident_t *Loc,
118118
for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
119119
SourceInfo Info(HostTargetMap.HstPtrName);
120120
INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
121-
DPxMOD " " DPxMOD " %-8" PRIuPTR " %-8" PRId64 " %s at %s:%d:%d\n",
121+
DPxMOD " " DPxMOD " %-8" PRIuPTR " %-8s %s at %s:%d:%d\n",
122122
DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
123123
HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin,
124-
HostTargetMap.getRefCount(), Info.getName(), Info.getFilename(),
125-
Info.getLine(), Info.getColumn());
124+
HostTargetMap.refCountToStr().c_str(), Info.getName(),
125+
Info.getFilename(), Info.getLine(), Info.getColumn());
126126
}
127127
Device.DataMapMtx.unlock();
128128
}

‎openmp/libomptarget/test/offloading/info.c‎

Lines changed: 10 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,10 @@
66

77
#define N 64
88

9+
#pragma omp declare target
10+
int global;
11+
#pragma omp end declare target
12+
913
extern void __tgt_set_info_flag(unsigned);
1014

1115
int main() {
@@ -19,10 +23,10 @@ int main() {
1923
// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
2024
// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
2125
// INFO: Libomptarget device 0 info: to(C[0:64])[256]
22-
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
23-
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
26+
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=A[0:64]
27+
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=B[0:64]
2428
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
25-
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
29+
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=C[0:64]
2630
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
2731
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
2832
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
@@ -45,6 +49,9 @@ int main() {
4549
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
4650
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
4751
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
52+
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
53+
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
54+
// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF unknown at unknown:0:0
4855
#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
4956
#pragma omp target firstprivate(val)
5057
{ val = 1; }

0 commit comments

Comments
 (0)