[OPENMP]Make __kmpc_push_tripcount thread safe.

[ABataev]

Summary

__kmpc_push_tripcount function is not thread safe and may lead to data
race when the target regions are executed in parallel threads. The patch
makes loopTripCnt counter thread aware and stores the tripcount value
per thread in the map. Access to map is guarded by mutex to prevent
data race in the map itself.
Test is for NVPTX target because it does not work correctly on the
host. Seems to me, there is a problem in libomp with target regions in
the parallel threads.

[jdoerfert]

This looks generally good to me and I think we should get this fix in. I have one nit inlined, and two follow up questions below:

1. Do we need to be interoperable with std::thread? I mean, would it be sufficient to ask for the OpenMP thread ID instead of the std::this_thread::get_id()?
2. Shouldn't we move the whole loop trip count logic into the target offload call? Afaik, we currently emit sth. like:

__kmpc_set_upcoming_trip_count(N);
__kmpc_target(....);

Now if we would move the N as a parameter into the "target" call we would not need to store it anyway, or am I missing something here?

libomptarget/src/omptarget.cpp

738 ↗

(On Diff #207559)

Nit: std::swap(ltc, I->second) and no braces.

[ABataev]

In D64080#1567062, @jdoerfert wrote:

This looks generally good to me and I think we should get this fix in. I have one nit inlined, and two follow up questions below:

1. Do we need to be interoperable with std::thread? I mean, would it be sufficient to ask for the OpenMP thread ID instead of the std::this_thread::get_id()?

It would be good to use OpenMP thread id, but, unfortunately, OpenMP thread id is not passed to the offloading functions (tgt_target etc.). That was the reason to use std::this_thread::get_id() instead of OpenMP thread id.
We can call __kmpc_global_thread_num(nullptr) directly instead. libomptarget already depends on libomp to support async offloading, so, maybe, we can use this libomp function instead of std::thread.

2. Shouldn't we move the whole loop trip count logic into the target offload call? Afaik, we currently emit sth. like:

__kmpc_set_upcoming_trip_count(N);
__kmpc_target(....);

Now if we would move the N as a parameter into the "target" call we would not need to store it anyway, or am I missing something here?

I thought about it and I think it is a good idea. But it is a completely different problem. We can't change the interface of the existing functions (for compatibility), instead, we need to provide the whole bunch of new functions with some new params. I just want to fix the problem with the existing kmpc_push_tripcount function. Redesign requires some additional work.

[ABataev]

1. Used std::swap.
2. Removed dependency on std::thread, used OpenMP thread id instead.

[jdoerfert]

In D64080#1567088, @ABataev wrote:

In D64080#1567062, @jdoerfert wrote:

This looks generally good to me and I think we should get this fix in. I have one nit inlined, and two follow up questions below:

1. Do we need to be interoperable with std::thread? I mean, would it be sufficient to ask for the OpenMP thread ID instead of the std::this_thread::get_id()?

It would be good to use OpenMP thread id, but, unfortunately, OpenMP thread id is not passed to the offloading functions (tgt_target etc.). That was the reason to use std::this_thread::get_id() instead of OpenMP thread id.
We can call __kmpc_global_thread_num(nullptr) directly instead. libomptarget already depends on libomp to support async offloading, so, maybe, we can use this libomp function instead of std::thread.

Looks better now.

2. Shouldn't we move the whole loop trip count logic into the target offload call? Afaik, we currently emit sth. like:

__kmpc_set_upcoming_trip_count(N);
__kmpc_target(....);

Now if we would move the N as a parameter into the "target" call we would not need to store it anyway, or am I missing something here?

I thought about it and I think it is a good idea. But it is a completely different problem. We can't change the interface of the existing functions (for compatibility), instead, we need to provide the whole bunch of new functions with some new params. I just want to fix the problem with the existing kmpc_push_tripcount function. Redesign requires some additional work.

Fair point. Let's put that on the TODO list and talk about what else we want to include on the mailing list.

Good to commit.

[grokos]

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

[ABataev]

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.
Plus, libomptarget already communicates with libomp using kmpc_omp_taskwait,though incorrectly. It uses 0 as thread ID always, though it must use real OpenMP thread ID, not 0. I'm going to fix this in the next patch.

[hfinkel]

In D64080#1567348, @ABataev wrote:

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.

If this should really be fixed in libomp, then we should fix it there. At the very least, there should be a FIXME about removing the workaround once the libomp fix is in place. @AndreyChurbanov, any thoughts on this?

[ABataev]

In D64080#1567533, @hfinkel wrote:

In D64080#1567348, @ABataev wrote:

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.

If this should really be fixed in libomp, then we should fix it there. At the very least, there should be a FIXME about removing the workaround once the libomp fix is in place. @AndreyChurbanov, any thoughts on this?

Hal, currently target tripcount has nothing to do with libomp. It is stored as part of the device in libomptarget. So, I think, to fix this problem in libomptarget is fine.
To put it into target task allocated memory is completely different problem, which requires some redesign and full support of target tasks. Plus, just like I said, it is not true that libomptarget is thread agnostic. Actually, it is not, since it uses __kmpc_omp_taskwait(NULL, 0) for dependendent target regions. Moreover, it is not thread safe too, because uses not real thread ID, but hardcoded 0 thread ID. It may lead to problems even in a single-threaded environment if the caller thread is not a master thread.

[hfinkel]

In D64080#1567632, @ABataev wrote:

In D64080#1567533, @hfinkel wrote:

In D64080#1567348, @ABataev wrote:

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.

If this should really be fixed in libomp, then we should fix it there. At the very least, there should be a FIXME about removing the workaround once the libomp fix is in place. @AndreyChurbanov, any thoughts on this?

Hal, currently target tripcount has nothing to do with libomp. It is stored as part of the device in libomptarget. So, I think, to fix this problem in libomptarget is fine.
To put it into target task allocated memory is completely different problem, which requires some redesign and full support of target tasks.

I think that it makes sense to note this in the code. Can you please add a comment along the lines of (assuming that I'm understanding you correctly):

// NOTE: Once libomp gains full target-task support, this state should be moved into the target task in libomp.

Plus, just like I said, it is not true that libomptarget is thread agnostic. Actually, it is not, since it uses __kmpc_omp_taskwait(NULL, 0) for dependendent target regions. Moreover, it is not thread safe too, because uses not real thread ID, but hardcoded 0 thread ID. It may lead to problems even in a single-threaded environment if the caller thread is not a master thread.

This is the issue to which "I'm going to fix this in the next patch." refers, correct?

[ABataev]

In D64080#1567797, @hfinkel wrote:

In D64080#1567632, @ABataev wrote:

In D64080#1567533, @hfinkel wrote:

In D64080#1567348, @ABataev wrote:

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.

If this should really be fixed in libomp, then we should fix it there. At the very least, there should be a FIXME about removing the workaround once the libomp fix is in place. @AndreyChurbanov, any thoughts on this?

Hal, currently target tripcount has nothing to do with libomp. It is stored as part of the device in libomptarget. So, I think, to fix this problem in libomptarget is fine.
To put it into target task allocated memory is completely different problem, which requires some redesign and full support of target tasks.

I think that it makes sense to note this in the code. Can you please add a comment along the lines of (assuming that I'm understanding you correctly):

// NOTE: Once libomp gains full target-task support, this state should be moved into the target task in libomp.

Ok.

Plus, just like I said, it is not true that libomptarget is thread agnostic. Actually, it is not, since it uses __kmpc_omp_taskwait(NULL, 0) for dependendent target regions. Moreover, it is not thread safe too, because uses not real thread ID, but hardcoded 0 thread ID. It may lead to problems even in a single-threaded environment if the caller thread is not a master thread.

This is the issue to which "I'm going to fix this in the next patch." refers, correct?

Yes.

[hfinkel]

In D64080#1567802, @ABataev wrote:

In D64080#1567797, @hfinkel wrote:

In D64080#1567632, @ABataev wrote:

In D64080#1567533, @hfinkel wrote:

In D64080#1567348, @ABataev wrote:

In D64080#1567327, @grokos wrote:

A long time ago we had identified the problem with the loop trip count and if I recall correctly the proposed solution was to store the trip count per OpenMP task. The fix would be implemented in libomp because libomptarget has no notion of tasks.

@AlexEichenberger Can you confirm whether or not this is the case? I may confuse this with something else...

The patch does exactly what you said: stores tripcount per task, but in the libomptarget, not libomp. When it will be implemented in libomp, we could remove this code. But we need fully functional implementation now, not in the future.

If this should really be fixed in libomp, then we should fix it there. At the very least, there should be a FIXME about removing the workaround once the libomp fix is in place. @AndreyChurbanov, any thoughts on this?

Hal, currently target tripcount has nothing to do with libomp. It is stored as part of the device in libomptarget. So, I think, to fix this problem in libomptarget is fine.
To put it into target task allocated memory is completely different problem, which requires some redesign and full support of target tasks.

I think that it makes sense to note this in the code. Can you please add a comment along the lines of (assuming that I'm understanding you correctly):

// NOTE: Once libomp gains full target-task support, this state should be moved into the target task in libomp.

Ok.

Thanks!

Plus, just like I said, it is not true that libomptarget is thread agnostic. Actually, it is not, since it uses __kmpc_omp_taskwait(NULL, 0) for dependendent target regions. Moreover, it is not thread safe too, because uses not real thread ID, but hardcoded 0 thread ID. It may lead to problems even in a single-threaded environment if the caller thread is not a master thread.

This is the issue to which "I'm going to fix this in the next patch." refers, correct?

Yes.

[ABataev]

Added note.

[hfinkel]

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

[jdoerfert]

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

It's one map per device, each map contains one trip count per OpenMP thread.

[hfinkel]

In D64080#1569283, @jdoerfert wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

It's one map per device, each map contains one trip count per OpenMP thread.

Ah, indeed. Nevermind.

[ABataev]

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

[hfinkel]

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

[ABataev]

In D64080#1569319, @hfinkel wrote:

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

Hard to tell, but definitely ee increase dependency between libomptarget and libomp. and, seems to me, threadprivate interface also uses some kind of hash table. We can use unordered_map instead.

[hfinkel]

In D64080#1569334, @ABataev wrote:

In D64080#1569319, @hfinkel wrote:

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

Hard to tell, but definitely ee increase dependency between libomptarget and libomp. and, seems to me, threadprivate interface also uses some kind of hash table. We can use unordered_map instead.

Given that we expect this to be temporary, I'm okay with it either way. unordered_map is probably a little bit better than std::map. I expect the number of entries in the map to be small because __kmpc_global_thread_num returns OpenMP thread ids, not system TIDs, right?

[ABataev]

In D64080#1569363, @hfinkel wrote:

In D64080#1569334, @ABataev wrote:

In D64080#1569319, @hfinkel wrote:

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

Hard to tell, but definitely ee increase dependency between libomptarget and libomp. and, seems to me, threadprivate interface also uses some kind of hash table. We can use unordered_map instead.

Given that we expect this to be temporary, I'm okay with it either way. unordered_map is probably a little bit better than std::map. I expect the number of entries in the map to be small because __kmpc_global_thread_num returns OpenMP thread ids, not system TIDs, right?

Yes.

[hfinkel]

In D64080#1569375, @ABataev wrote:

In D64080#1569363, @hfinkel wrote:

In D64080#1569334, @ABataev wrote:

In D64080#1569319, @hfinkel wrote:

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

Hard to tell, but definitely ee increase dependency between libomptarget and libomp. and, seems to me, threadprivate interface also uses some kind of hash table. We can use unordered_map instead.

Given that we expect this to be temporary, I'm okay with it either way. unordered_map is probably a little bit better than std::map. I expect the number of entries in the map to be small because __kmpc_global_thread_num returns OpenMP thread ids, not system TIDs, right?

Yes.

SGTM.

[hfinkel]

In D64080#1569501, @hfinkel wrote:

In D64080#1569375, @ABataev wrote:

In D64080#1569363, @hfinkel wrote:

In D64080#1569334, @ABataev wrote:

In D64080#1569319, @hfinkel wrote:

In D64080#1569299, @ABataev wrote:

In D64080#1569266, @hfinkel wrote:

In D64080#1568348, @ABataev wrote:

Added note.

The patch description says that you add a trip count "per thread", but the code is adding a trip count per device. Could this still cause a problem if multiple threads have target regions for the same device?

No. We already have tripcount per device: each device stores its unique tripcount value, but only one. I extended the existing solution to have the tripcount not only per device, but also per host thread.

Would there be any advantage to using __kmpc_threadprivate_register and __kmpc_threadprivate (or similar) instead of storing a map with thread ids?

Hard to tell, but definitely ee increase dependency between libomptarget and libomp. and, seems to me, threadprivate interface also uses some kind of hash table. We can use unordered_map instead.

Given that we expect this to be temporary, I'm okay with it either way. unordered_map is probably a little bit better than std::map. I expect the number of entries in the map to be small because __kmpc_global_thread_num returns OpenMP thread ids, not system TIDs, right?

Yes.

SGTM.

A note on naming, but this also LGTM.

I chatted offline with Johannes, and I agree with him, that longer term we should have a new interface that doesn't require this "push" precall. Regardless, I think that the comment is fine, in that, so long as we have to support this interface (which we might need for backward compatibility even after we add a new interface), it's fine to note where a cleaner location for that state data might be.

libomptarget/src/device.h

101 ↗

(On Diff #207778)

While we're changing this, we should also update it to following our naming conventions:

loopTripCnt -> LoopTripCnt

[ABataev]

libomptarget/src/device.h

101 ↗

(On Diff #207778)

I thin the renaming better to put in another patch, but if you're ok with this, I'll rename it here

[hfinkel]

libomptarget/src/device.h

101 ↗

(On Diff #207778)

My thinking is: Because we need to change every line of code which references it anyway, changing the name along with changing the type makes sense. To me, I don't think it matters whether you split that into one patch or two. If you think it's cleaner to changing the name in a separate patch from changing the type, that's fine too.

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/tripcount.c

+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int res = 0;
+
+#pragma omp parallel num_threads(2) reduction(+:res)
+  {
+    int tid = omp_get_thread_num();
+#pragma omp target teams distribute reduction(+:res)
+    for (int i = tid; i < 2; i++)
+      ++res;
+  }
+  // The first thread makes 2 iterations, the second - 1. Expected result of the
+  // reduction res is 3.
+
+  // CHECK: res = 3.
+  printf("res = %d.\n", res);
+  return 0;
+}

openmp/trunk/libomptarget/src/device.h

struct DeviceTy {
 
   HostDataToTargetListTy HostDataToTargetMap;
   PendingCtorsDtorsPerLibrary PendingCtorsDtors;
 
   ShadowPtrListTy ShadowPtrMap;
 
   std::mutex DataMapMtx, PendingGlobalsMtx, ShadowMtx;
 
-  uint64_t loopTripCnt;
+  // NOTE: Once libomp gains full target-task support, this state should be
+  // moved into the target task in libomp.
+  std::map<int32_t, uint64_t> loopTripCnt;
 
   int64_t RTLRequiresFlags;
 
   DeviceTy(RTLInfoTy *RTL)
       : DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
         HasPendingGlobals(false), HostDataToTargetMap(),
         PendingCtorsDtors(), ShadowPtrMap(), DataMapMtx(), PendingGlobalsMtx(),
-        ShadowMtx(), loopTripCnt(0), RTLRequiresFlags(0) {}
+        ShadowMtx(), RTLRequiresFlags(0) {}
 
   // The existence of mutexes makes DeviceTy non-copyable. We need to
   // provide a copy constructor and an assignment operator explicitly.
   DeviceTy(const DeviceTy &d)
       : DeviceID(d.DeviceID), RTL(d.RTL), RTLDeviceID(d.RTLDeviceID),
         IsInit(d.IsInit), InitFlag(), HasPendingGlobals(d.HasPendingGlobals),
         HostDataToTargetMap(d.HostDataToTargetMap),
         PendingCtorsDtors(d.PendingCtorsDtors), ShadowPtrMap(d.ShadowPtrMap),

openmp/trunk/libomptarget/src/interface.cpp

EXTERN int __tgt_target_teams_nowait(int64_t device_id, void *host_ptr,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
     __kmpc_omp_taskwait(NULL, 0);
 
   return __tgt_target_teams(device_id, host_ptr, arg_num, args_base, args,
                             arg_sizes, arg_types, team_num, thread_limit);
 }
 
-
-// The trip count mechanism will be revised - this scheme is not thread-safe.
 EXTERN void __kmpc_push_target_tripcount(int64_t device_id,
     uint64_t loop_tripcount) {
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     DP("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false);
     return;
   }
 
   DP("__kmpc_push_target_tripcount(%" PRId64 ", %" PRIu64 ")\n", device_id,
       loop_tripcount);
-  Devices[device_id].loopTripCnt = loop_tripcount;
+  TblMapMtx.lock();
+  Devices[device_id].loopTripCnt.emplace(__kmpc_global_thread_num(NULL),
+                                         loop_tripcount);
+  TblMapMtx.unlock();
 }

openmp/trunk/libomptarget/src/omptarget.cpp

#endif
     tgt_args.push_back(TgtPtrBegin);
     tgt_offsets.push_back(TgtBaseOffset);
   }
 
   assert(tgt_args.size() == tgt_offsets.size() &&
       "Size mismatch in arguments and offsets");
 
   // Pop loop trip count
-  uint64_t ltc = Device.loopTripCnt;
-  Device.loopTripCnt = 0;
+  uint64_t ltc = 0;
+  TblMapMtx.lock();
+  auto I = Device.loopTripCnt.find(__kmpc_global_thread_num(NULL));
+  if (I != Device.loopTripCnt.end())
+    std::swap(ltc, I->second);
+  TblMapMtx.unlock();
 
   // Launch device execution.
   DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
       TargetTable->EntriesBegin[TM->Index].name,
       DPxPTR(TargetTable->EntriesBegin[TM->Index].addr), TM->Index);
   if (IsTeamConstruct) {
     rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr,
         &tgt_args[0], &tgt_offsets[0], tgt_args.size(), team_num,

openmp/trunk/libomptarget/src/private.h

 
 // Implemented in libomp, they are called from within __tgt_* functions.
 #ifdef __cplusplus
 extern "C" {
 #endif
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_omp_taskwait(void *loc_ref, int32_t gtid) __attribute__((weak));
+int32_t __kmpc_global_thread_num(void *) __attribute__((weak));
 int __kmpc_get_target_offload(void) __attribute__((weak));
 #ifdef __cplusplus
 }
 #endif
 
 #ifdef OMPTARGET_DEBUG
 extern int DebugLevel;