Automatic asynchronous execution of OpenMP Target Regions

[randreshg]

Summary

This patch allows Automatic asynchronous execution of OpenMP Target Regions.
When the environment variable LIBOMPTARGET_INTRA_THREAD_ASYNC is enabled
(LIBOMPTARGET_INTRA_THREAD_ASYNC=1), the implicit barrier that exists at the end
of every target region is removed and synchronization is only performed on memory transfers
or at the end of the program when the OpenMP runtime calls its destructors

[ye-luo]

Could you please explain why removing synchronization at the end of a target region (without nowait) is a valid optimization. Also there should be no valid program relying on synchronization at the end of the program"

[josemonsalve2]

Ye,

It is only valid under the assumption of non unified shared memory. State across host and device is only visible during data movements. So it is up until then when changes in the host or device data is reflected. Assuming there are no external runtimes, it is possible to synchronize only on data movements, conserving the data dependencies between host and device.

[ye-luo]

When I ran target region A and then B and B consumes numbers generated by A on the device and there is no transfer involved. B may start before the A get the numbers ready.

[randreshg]

As long as the RTL of the device provides a queue mechanism to execute target tasks sequentially, the dependencies among different tasks will be respected, for NVIDIA GPUs the queue mechanism is the stream. This patch, in contrast with the current OpenMP offloading implementation, launches the execution of tasks into the same stream.

[josemonsalve2]

CUDA streams are FIFO queues. But it’s true this will not work if the device queue is not FIFO. In the case of CUDA this works for the case you described without being a read after write dependency.

[ye-luo]

1. right now in CUDA plugin, streams are pooled. There is no guarantee that A and B get the same Stream when multiple threads all are doing their own As and Bs.
2. Need to minimize designing libomptarget based on CUDA behaviors.

[josemonsalve2]

I agree on 2. Any recommendations? We can move some of the logic there.

Regarding 1 this patch does not return the stream to the pool right away. It is held by the thread until a synchronization occurs. Synchronizations aré evaluated lazily on data movement. But the stream does not change between consecutive target invokes with no data movement

[ye-luo]

In D132306#3737736, @josemonsalve2 wrote:

I agree on 2. Any recommendations? We can move some of the logic there.

Regarding 1 this patch does not return the stream to the pool right away. It is held by the thread until a synchronization occurs. Synchronizations aré evaluated lazily on data movement. But the stream does not change between consecutive target invokes with no data movement

Explicitly tied to thread context is even worse. If I wrap both region with openmp CPU tasks and make task B depends on task A. task A completes but the kernel is still flying. Task B may start at anytime and there is no guarantee of running on the same thread.

[josemonsalve2]

That one is a good point. Let us revise that. Since the effect we actually want to have is the creation of the task graph sequentially.

[jdoerfert]

openmp/libomptarget/include/device.h

340

Why is this static, that seems wrong.

475

Documentation, please.

openmp/libomptarget/include/omptarget.h

198–203

You cannot just remove this. See https://reviews.llvm.org/D132045, as it introduces a flag you can use to disable synchronization here.

openmp/libomptarget/src/device.cpp

69

Reverse the condition, single line alternative comes first, then the complex consequence (without the else).

81–83

Use early exits and no else after return.

[josemonsalve2]

openmp/libomptarget/src/device.cpp

61–65

It may be better to merge these into:

DP("Asynchronous execution %s\n", AsyncFlag ? "Enabled" : "Disabled");

[randreshg]

openmp/libomptarget/include/device.h

340

The idea of AsyncInfoMng is to have a way to control the AsyncInfo object to skip synchronization that is not needed. It's part of the Device class, so we can sync like this:

Device.syncAsyncInfo(AsyncInfo, true);

however, it is not necessary for every device to have a copy of AsyncInfoMng. That's why it's static.

[randreshg]

diff updated.

• The AsyncInfoManager is now part of the omptarget.h and not device.h
• Global variable AIM was added in interface.cpp
• Flag was added to the Synchronization in asyncinfo destructor

[randreshg]

Documentation added

[jdoerfert]

Some more initial comments

openmp/docs/optimizations/OpenMPOpt.rst

113

Doesn't the length need to match?

116

For nowait regions this is maybe misleading. We should probably say they can be executed synchronously depending on the pragmas and implementation. In that case ...

134

This is misleading. Rather explain what would happen if the map on the target was triggering a transfer.

openmp/libomptarget/include/omptarget.h

232

This default is dangerous. Either swap it or avoid the default. Also, explain what it means to synchronize but not to force synchronize.

236

Use doxygen comments /// not //.

openmp/libomptarget/src/device.cpp

52

Unrelated

openmp/libomptarget/src/interface.cpp

25

We should not have a global like this. It probably belongs in/to the Device object.

254–259

A comment would be nice as this is the one place we do not force synchronization.

[randreshg]

This new version has the following changes:

• Updated documentation
• The AsyncInfoManager is part of the Device class and now has a map that contains thread ids as its key values.

[randreshg]

ping

[jdoerfert]

openmp/docs/optimizations/OpenMPOpt.rst

120

135–140

openmp/libomptarget/include/device.h

320

Documentation, please.

353

I would rather expose the AIM than have 3 functions that just forward to it.

openmp/libomptarget/include/omptarget.h

26

We can use LLVM data structures now too, e.g., maps.

openmp/libomptarget/src/interface.cpp

254–259

multi-line conditionals should have braces. The comment is unhelpful. It states what the code already says, not why the code is this way.

openmp/libomptarget/src/omptarget.cpp

58

Lot's of static flags to simply lookup a env var once. Why don't we do it the same way as for other env vars?

60

66

67

You did get an iterator, why do another lookup via [...]? You can do a single [...] lookup and use the reference result to check and update it.

69

And this is the 3rd lookup into the map..

70

Do we ever have to clear the map?

87

I doubt we need both these lines.

[randreshg]

This new patch addresses all previous comments from reviewers

[jdoerfert]

I know this is getting tiresome but we need to make sure people understand what's happening and this plays well with future extensions. More comments.

openmp/libomptarget/include/omptarget.h

195

ShouldSyncWhenDestroyed

232

You never return a nullptr, make it a reference.

240

The description is unhelpful. What AsyncInfo, etc. The function is also unused, do we need it?

openmp/libomptarget/src/interface.cpp

100

Now this pattern is somewhat unfortunate. You get the AsyncInfo from the AIM and then you need to be careful to call the right synchronize. If the AsynFlag was global you could move the new "sync" logic into the regular AsyncInfo sync, right? AIM would just be used to manage the map ID -> AsyncInfo. WDYT about this scheme? You could check the env variable once, like we do it for some others: https://github.com/llvm/llvm-project/blob/23bc343855fdf6fb7668abadf2b064034b207981/openmp/libomptarget/src/rtl.cpp#L43

openmp/libomptarget/src/omptarget.cpp

59–62

[randreshg]

openmp/libomptarget/src/interface.cpp

100

Agreed!
The pattern you suggest allows for separating the synchronization logic and AsyncInfo management.

[randreshg]

Ping

[jdoerfert]

I think this is fine for now. @ye-luo can this go in as opt-in feature that we probably refine as we go?

openmp/libomptarget/src/interface.cpp

100

Better. Not super happy about the explicit free call but that's fine for now.

[ye-luo]

How many testing has been done? it seems only workable on toy examples.

openmp/docs/optimizations/OpenMPOpt.rst

140

Through this example, I only see TT1 and TT2 racing when the async feature is enabled.

openmp/libomptarget/src/device.cpp

55

Why AIM captures a pointer instead of reference?

openmp/libomptarget/src/omptarget.cpp

30–38

It is mess. synchroning or not depends on a bunch of states. Move the if-statement to the caller side and make the logic directly exposed on the use side.

76

I think this is against coding principles, pass in a reference and delete its memory.

openmp/libomptarget/src/rtl.cpp

42

Is AsyncFlag documented?

[ye-luo]

openmp/libomptarget/include/omptarget.h

195

Add const

220

This is insufficient documentation. Please explain what this struct does actually not just what it is used for.

221

Why is this a struct instead of a class? It encourage code like AIM.AsyncInfoM[1]? Use class and appropriate private/public.

222

Why AsyncInfoTy needs to be associated with thread::id. Does this add unnecessary entanglement between target tasks.

[ye-luo]

openmp/libomptarget/include/omptarget.h

222

it is necessary to document the design choice that AsyncInfoTy objects and threadIDs have one to one mapping when AsyncFlag is true.

openmp/libomptarget/src/omptarget.cpp

59

New a pointer and then return its reference. "reference" types should not be used to manage the ownership.

[tianshilei1992]

openmp/libomptarget/src/rtl.cpp

42

I don't think this is a good design, that tries to modify the state of libomptarget from plugins, especially the corresponding env is called LIBOMPTARGET_ASYNC, which indicates it should be handled in libomptarget instead of in each plugin. Potentially I think we could add a plugin interface function to tell what opt-in feature is enabled. This could be done in a separate patch, but is better to land it before this patch. I don't think it makes sense to open the door (now) and then close it later, because I don't think "later" will come.

BTW, LIBOMPTARGET_ASYNC is a little bit confusing. It's too general.

[tianshilei1992]

openmp/libomptarget/include/omptarget.h

21–22

It is recommended to add a blank line between LLVM and STL headers.

222

Better to use LLVM ADT here

[tianshilei1992]

openmp/libomptarget/src/rtl.cpp

42

nvm, I didn't look at it right. It's in libomptarget. My bad. But the env name is true.

[jdoerfert]

In D132306#3886372, @ye-luo wrote:

How many testing has been done? it seems only workable on toy examples.

Wrt toy examples:

openmp/docs/optimizations/OpenMPOpt.rst

140

As discussed yesterday, there is no race.

openmp/libomptarget/include/omptarget.h

222

This does not depend on tasks but threads. If threads are independent wrt offload, this extension allows them to run host and device code asynchronously. If threads are not independent wrt. offload, this extension cannot be used.

openmp/libomptarget/src/omptarget.cpp

30–38

Exposing this to the user side is not only not helpful but will actively harm things. The logic depends on two flags, not "a bunch of states". One is passed by the user, one describes the system setup. This is totally fine.

59

What about this:

Replace "get" with

AsyncInfoMng::register(AsyncInfo &AI);

Remove "free".
Register is implemented as no-op w/o the AsyncInfo flag set.
Otherwise it'll replace AI with a dynamically allocated one.

openmp/libomptarget/src/rtl.cpp

42

What about LIBOMPTARGET_INTRA_THREAD_ASYNC?

[tianshilei1992]

openmp/libomptarget/src/rtl.cpp

42

that sounds good!

[ye-luo]

openmp/libomptarget/include/omptarget.h

222

This does not depend on tasks but threads. If threads are independent wrt offload, this extension allows them to run host and device code asynchronously. If threads are not independent wrt. offload, this extension cannot be used.

Fair enough. This feature relies on thread id to impose dependency. It can only be used under certain restrictions.

openmp/libomptarget/src/omptarget.cpp

30–38

Two flags = 4 states. Not a low number of variants. Calling synchronize() but it may or may not do the sync. When should the user set ForceSync to true? Better to have some explanation.

59

Not getting what you mean.
I'm still expecting, AsyncInfoTy object being destroyed properly at the end of the target region, when AsyncInfo=false.

[jdoerfert]

openmp/libomptarget/src/omptarget.cpp

30–38

4 states, ok. However, they collapse to 2; it's a single conditional after all: synchronize or not.
The rules are fairly simple (one or condition) and documented:

/// Synchronize all pending actions when the LIBOMPTARGET_ASYNC env var
/// is disabled or when synchronization is forced (ForceSync = true) 
// Otherwise, synchronization is skipped
/// \returns OFFLOAD_FAIL or OFFLOAD_SUCCESS appropriately.

59

I'm still expecting, AsyncInfoTy object being destroyed properly at the end of the target region, when AsyncInfo=false.

It is properly destroyed with this patch, and it will be with the proposed scheme. In the proposed scheme we have a local AsyncInfo (as we have upstream now) and we use it if AsyncInfo=false.

[randreshg]

openmp/libomptarget/src/omptarget.cpp

59

Im not following this. Could you please provide a pseudocode?

[randreshg]

This patch addresses comments from reviewers:

• Env var name is LIBOMPTARGET_INTRA_THREAD_ASYNC
• The pattern to get and destroy the AsyncInfoTy object changed.
• A DenseMap is used instead of the std::map.

[randreshg]

ping

[randreshg]

ping

[randreshg]

Patch updated to the trunk version.
Changes compared to the last patch:

• the HasDataTransfer flag was added.
• the functions AsyncInfoTy *get() was added to both AsyncInfoTy and TaskAsyncInfoWrapperTy

[josemonsalve2]

Adding another limitation to this approach. In the following code:

void aaaa(int b) {
    // TARGET A
    #pragma omp target
    {}

    // TARGET B
    #pragma omp target nowait
    {}

    // TARGET C
    #pragma omp target
    {}
}

The task A->B dependency should be respected, since A is originally synchronous, it should have been executed before B. However, B will be spawned in a different thread and A can potentially execute after B.

openmp/libomptarget/include/omptarget.h

208

It may be a good idea to add a comment here explaining why this is necessary. It has to do with garbage collection of the device RTL, as per our conversation today.

[RaviNarayanaswamy]

In D132306#4207335, @josemonsalve2 wrote:

Adding another limitation to this approach. In the following code:

void aaaa(int b) {
    // TARGET A
    #pragma omp target
    {}

    // TARGET B
    #pragma omp target nowait
    {}

    // TARGET C
    #pragma omp target
    {}
}

The task A->B dependency should be respected, since A is originally synchronous, it should have been executed before B. However, B will be spawned in a different thread and A can potentially execute after B.

The thread which encounter A will not reach B until A has completed so how can B be spawned when A is active.

[josemonsalve2]

Hi Ravi,

The purpose of this is to have regions with nowait to be asynchronous (but ordered within the stream). (See paper). So using the solution presented by @randreshg will have this issue. The problem is that the encountering thread will push A to the queue, but B is going to be lowered as a host task. B can be pushed into another thread, and potentially delayed way before A finishes. A must be synchronized such that the original order is maintained. My example was rather simplistic, but it demonstrates the issue. While A and C will be in order, B will not be ordered w.r.t. A.

Of course this is not default OpenMP behavior where what you said does apply.

In D132306#4207789, @RaviNarayanaswamy wrote:

The thread which encounter A will not reach B until A has completed so how can B be spawned when A is active.

[randreshg]

Test added

openmp/docs/optimizations/OpenMPOpt.rst

 
 These optimizations can have very large performance implications. Both of these
 optimizations rely heavily on inter-procedural analysis. Because of this,
 offloading applications should ideally be contained in a single translation unit
 and functions should not be externally visible unless needed. OpenMPOpt will
 inform the user if any globalization calls remain if remarks are enabled. This
 should be treated as a defect in the program.
 
+.. _Others:
+
+Others
+=========
+
+.. contents::
+   :local:
+   :depth: 1
+
+Automatic Asynchronous Execution of Target Regions
+--------------------------------------------------
+
+Offloaded regions can be executed synchronously or asynchronously
+depending on the pragmas instruction (e.g. nowait) and implementation.
+The environment flag  `LIBOMPTARGET_INTRA_THREAD_ASYNC=1` allows asynchronous 
+execution of OpenMP Target Regions on Nvidia GPUs by synchronizing 
+only on memory transfers or at the end of the program when the 
+OpenMP runtime is torn down.
+
+.. code-block:: c++
+  LIBOMPTARGET_ASYNC=1 ./simple1 //For async execution
+
+For example:
+.. code-block:: c++
+  #pragma omp target enter data map(alloc: array[:N])
+  #pragma omp target map(tofrom: array[:N])
+  busy_device(array); //TT1
+  #pragma omp target map(tofrom: array[:N])
+  busy_device(array); //TT2
+  busy_host();
+  #pragma omp target exit data map(from: array[:N])
+
+In the example above, the target regions inside the enter/exit data region
+do not perform any memory transfer, since the mapped memory is already
+on the device. So, the host launches the execution of both target 
+regions (TT1 and TT2) in the device queue and continues the execution
+on its end (busy host). Offloaded regions that trigger memory transfers 
+will not benefit from this optimization.
+
+
+Limitations:
+- Cross-device synchronization primitives or atomics are not allowed.
+- Timers wrapping target regions will not show the time it takes to
+  execute the target task but the time the runtime takes to launch
+  the execution of the task.
+- Synchronization using different threads: Target task and
+  synchronizations in different threads is not supported yet
+  when the flag for asynchronous asynchronous execution
+  is enabled.
+- RTL of the target device should allow asynchronous execution and 
+  provide a queue mechanism to execute kernels sequentially.
+
 Resources
 =========
 
 - 2021 OpenMP Webinar: "A Compiler's View of OpenMP" https://youtu.be/eIMpgez61r4
 - 2020 LLVM Developers’ Meeting: "(OpenMP) Parallelism-Aware Optimizations" https://youtu.be/gtxWkeLCxmU
 - 2019 EuroLLVM Developers’ Meeting: "Compiler Optimizations for (OpenMP) Target Offloading to GPUs" https://youtu.be/3AbS82C3X30

openmp/libomptarget/include/device.h

 ///
 struct PendingCtorDtorListsTy {
   std::list<void *> PendingCtors;
   std::list<void *> PendingDtors;
 };
 typedef std::map<__tgt_bin_desc *, PendingCtorDtorListsTy>
     PendingCtorsDtorsPerLibrary;
 
+class AsyncInfoMng;
+
 struct DeviceTy {
   int32_t DeviceID;
   RTLInfoTy *RTL;
   int32_t RTLDeviceID;
 
   bool IsInit;
   std::once_flag InitFlag;
   bool HasPendingGlobals;
 
   /// Host data to device map type with a wrapper key indirection that allows
   /// concurrent modification of the entries without invalidating the underlying
   /// entries.
   using HostDataToTargetListTy =
       std::set<HostDataToTargetMapKeyTy, std::less<>>;
 
   /// The HDTTMap is a protected object that can only be accessed by one thread
   /// at a time.
   ProtectedObj<HostDataToTargetListTy> HostDataToTargetMap;
 
   /// The type used to access the HDTT map.
   using HDTTMapAccessorTy = decltype(HostDataToTargetMap)::AccessorTy;
 
   PendingCtorsDtorsPerLibrary PendingCtorsDtors;
 
   ShadowPtrListTy ShadowPtrMap;
 
   std::mutex PendingGlobalsMtx, ShadowMtx;
 
+  /// Manages AsyncInfo associated to the device and thread
+  /// depending on the value of the env variable LIBOMPTARGET_INTRA_THREAD_ASYNC
+  AsyncInfoMng AIM;
+
   DeviceTy(RTLInfoTy *RTL);
   // DeviceTy is not copyable
   DeviceTy(const DeviceTy &D) = delete;
   DeviceTy &operator=(const DeviceTy &D) = delete;
 
   ~DeviceTy();
 
   // Return true if data can be copied to DstDevice directly
   bool isDataExchangable(const DeviceTy &DstDevice);
 
   /// Lookup the mapping of \p HstPtrBegin in \p HDTTMap. The accessor ensures
   /// exclusive access to the HDTT map.
   LookupResult lookupMapping(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
                              int64_t Size);
 
   /// Get the target pointer based on host pointer begin and base. If the
   /// mapping already exists, the target pointer will be returned directly. In
   /// addition, if required, the memory region pointed by \p HstPtrBegin of size
   /// \p Size will also be transferred to the device. If the mapping doesn't
struct DeviceTy {
 
   /// Deallocate \p LR and remove the entry. Assume the total reference count is
   /// zero and the calling thread is the deleting thread for \p LR. \p HDTTMap
   /// ensure the caller holds exclusive access and can modify the map. Return \c
   /// OFFLOAD_SUCCESS if the map entry existed, and return \c OFFLOAD_FAIL if
   /// not. It is the caller's responsibility to skip calling this function if
   /// the map entry is not expected to exist because \p HstPtrBegin uses shared
   /// memory.
-  int deallocTgtPtr(HDTTMapAccessorTy &HDTTMap, LookupResult LR, int64_t Size);
+  int deallocTgtPtr(HDTTMapAccessorTy &HDTTMap, LookupResult LR, int64_t Size,
+                    AsyncInfoTy &AsyncInfo);
 
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL
   int32_t initOnce();
   __tgt_target_table *loadBinary(void *Img);
 

 private:
   // Call to RTL
   void init(); // To be called only via DeviceTy::initOnce()
 
   /// Deinitialize the device (and plugin).
   void deinit();
 };
 
 extern bool deviceIsReady(int DeviceNum);
 
 /// Struct for the data required to handle plugins
 struct PluginManager {
   PluginManager(bool UseEventsForAtomicTransfers)
       : UseEventsForAtomicTransfers(UseEventsForAtomicTransfers) {}
 
   /// RTLs identified on the host
   RTLsTy RTLs;

openmp/libomptarget/include/omptarget.h

 //===-------- omptarget.h - Target independent OpenMP target RTL -- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Interface to be used by Clang during the codegen of a
 // target region.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_H_
 #define _OMPTARGET_H_
 
+#include <SourceInfo.h>
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/DenseMapInfo.h"
+
 #include <deque>
+#include <mutex>
 #include <stddef.h>
 #include <stdint.h>
-
-#include <SourceInfo.h>
+#include <thread>
 
 #define OFFLOAD_SUCCESS (0)
 #define OFFLOAD_FAIL (~0)
 
 #define OFFLOAD_DEVICE_DEFAULT -1
 
 // Don't format out enums and structs.
 // clang-format off
 

 /// mistakes.
 class AsyncInfoTy {
   /// Locations we used in (potentially) asynchronous calls which should live
   /// as long as this AsyncInfoTy object.
   std::deque<void *> BufferLocations;
 
   __tgt_async_info AsyncInfo;
   DeviceTy &Device;
+  /// Enables/Disable synchronization in destructor
+  const bool ShouldSyncWhenDestroyed;
 
 public:
-  AsyncInfoTy(DeviceTy &Device) : Device(Device) {}
-  ~AsyncInfoTy() { synchronize(); }
+  AsyncInfoTy(DeviceTy &Device, bool ShouldSyncWhenDestroyed = true)
+      : Device(Device), ShouldSyncWhenDestroyed(ShouldSyncWhenDestroyed) {}
+  ~AsyncInfoTy() {
+    if (ShouldSyncWhenDestroyed)
+      synchronize(true);
+  }
 
   /// Implicit conversion to the __tgt_async_info which is used in the
   /// plugin interface.
   operator __tgt_async_info *() { return &AsyncInfo; }
 
-  /// Synchronize all pending actions.
-  ///
+  /// Synchronize all pending actions when the LIBOMPTARGET_INTRA_THREAD_ASYNC
+  /// env var is disabled or when synchronization is forced (ForceSync = true)
+  /// Otherwise, synchronization is skipped
   /// \returns OFFLOAD_FAIL or OFFLOAD_SUCCESS appropriately.
-  int synchronize();
+  int synchronize(bool ForceSync = false);
 
   /// Return a void* reference with a lifetime that is at least as long as this
   /// AsyncInfoTy object. The location can be used as intermediate buffer.
   void *&getVoidPtrLocation();
 };
 
+/// This structs allows for automatic asynchronous execution of target regions.
+/// It associates an AsyncInfoTy object with a thread id to guarantee that
+/// target tasks of the same thread are launched on the same queue.
+class AsyncInfoMng {
+  llvm::DenseMap<std::thread::id, AsyncInfoTy *> AsyncInfoM;
+  DeviceTy &Device;
+  std::mutex AsyncMtx;
+
+public:
+  AsyncInfoMng(DeviceTy &Device);
+  ~AsyncInfoMng();
+
+  /// Replaces AI with the associated AsyncInfo object when
+  /// LIBOMPTARGET_INTRA_THREAD_ASYNC is enabled.
+  AsyncInfoTy &registerAI(AsyncInfoTy *AI);
+};
+
+/// Specialize DenseMapInfo for std::thread::id
+template <> struct llvm::DenseMapInfo<std::thread::id> {
+
+  static std::thread::id getEmptyKey() {
+      return std::thread::id();
+  }
+
+  static std::thread::id getTombstoneKey() {
+    return std::thread::id();
+  }
+
+  static unsigned getHashValue(const std::thread::id &Val) {
+    static std::hash<std::thread::id> hasher;
+    return hasher(Val);
+  }
+
+  static bool isEqual(std::thread::id LHS, std::thread::id RHS) {
+    return LHS == RHS;
+  }
+};
+    
 /// This struct is a record of non-contiguous information
 struct __tgt_target_non_contig {
   uint64_t Offset;
   uint64_t Count;
   uint64_t Stride;
 };
 
 struct __tgt_device_info {

openmp/libomptarget/src/device.cpp

int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
   }
 
   if (NeedNewEvent)
     setEvent(Event);
 
   return OFFLOAD_SUCCESS;
 }
 
 DeviceTy::DeviceTy(RTLInfoTy *RTL)
     : DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
       HasPendingGlobals(false), PendingCtorsDtors(), ShadowPtrMap(),
-      PendingGlobalsMtx(), ShadowMtx() {}
+      PendingGlobalsMtx(), ShadowMtx(), AIM(*this) {}
 
 DeviceTy::~DeviceTy() {
   if (DeviceID == -1 || !(getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE))
     return;
 
   ident_t Loc = {0, 0, 0, 0, ";libomptarget;libomptarget;0;0;;"};
   dumpTargetPointerMappings(&Loc, *this);
 }
 
 int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
   HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
 
   // Check if entry exists
   auto It = HDTTMap->find(HstPtrBegin);
   if (It != HDTTMap->end()) {
     HostDataToTargetTy &HDTT = *It->HDTT;
     // Mapping already exists
     bool IsValid = HDTT.HstPtrEnd == (uintptr_t)HstPtrBegin + Size &&
                    HDTT.TgtPtrBegin == (uintptr_t)TgtPtrBegin;
     if (IsValid) {
       DP("Attempt to re-associate the same device ptr+offset with the same "
          "host ptr, nothing to do\n");
       return OFFLOAD_SUCCESS;
     }
     REPORT("Not allowed to re-associate a different device ptr+offset with "
            "the same host ptr\n");
     return OFFLOAD_FAIL;
   }
 
   // Mapping does not exist, allocate it with refCount=INF
   const HostDataToTargetTy &NewEntry =
       *HDTTMap
            ->emplace(new HostDataToTargetTy(
                /*HstPtrBase=*/(uintptr_t)HstPtrBegin,
                /*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
                /*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
if (LR.Flags.IsContained || LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) {
     uintptr_t TP = HT.TgtPtrBegin + (HP - HT.HstPtrBegin);
     return (void *)TP;
   }
 
   return NULL;
 }
 
 int DeviceTy::deallocTgtPtr(HDTTMapAccessorTy &HDTTMap, LookupResult LR,
-                            int64_t Size) {
+                            int64_t Size, AsyncInfoTy &AsyncInfo) {
   // Check if the pointer is contained in any sub-nodes.
   if (!(LR.Flags.IsContained || LR.Flags.ExtendsBefore ||
         LR.Flags.ExtendsAfter)) {
     REPORT("Section to delete (hst addr " DPxMOD ") does not exist in the"
            " allocated memory\n",
            DPxPTR(LR.Entry->HstPtrBegin));
     return OFFLOAD_FAIL;
   }
 
   auto &HT = *LR.Entry;
   // Verify this thread is still in charge of deleting the entry.
   assert(HT.getTotalRefCount() == 0 &&
          HT.getDeleteThreadId() == std::this_thread::get_id() &&
          "Trying to delete entry that is in use or owned by another thread.");
-
+  // We need to force synchronization before deallocating data.
+  int Ret = AsyncInfo.synchronize(true);
+  if (Ret != OFFLOAD_SUCCESS)
+    return OFFLOAD_FAIL;
+  // Delete tgt data
   DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
      DPxPTR(HT.TgtPtrBegin), Size);
   deleteData((void *)HT.TgtPtrBegin);
   INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
        "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
        ", Size=%" PRId64 ", Name=%s\n",
        DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size,
        (HT.HstPtrName) ? getNameFromMapping(HT.HstPtrName).c_str() : "unknown");
   void *Event = LR.Entry->getEvent();
   HDTTMap->erase(LR.Entry);
   delete LR.Entry;
 
-  int Ret = OFFLOAD_SUCCESS;
+  Ret = OFFLOAD_SUCCESS;
   if (Event && destroyEvent(Event) != OFFLOAD_SUCCESS) {
     REPORT("Failed to destroy event " DPxMOD "\n", DPxPTR(Event));
     Ret = OFFLOAD_FAIL;
   }
 
   return Ret;
 }
 

openmp/libomptarget/src/interface.cpp

 #include "rtl.h"
 
 #include <cassert>
 #include <cstdio>
 #include <cstdlib>
 #include <mutex>
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t Flags) {
   TIMESCOPE();
   PM->RTLs.registerRequires(Flags);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds a target shared library to the target execution image
 EXTERN void __tgt_register_lib(__tgt_bin_desc *Desc) {
for (int I = 0; I < ArgNum; ++I) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        I, DPxPTR(ArgsBase[I]), DPxPTR(Args[I]), ArgSizes[I], ArgTypes[I],
        (ArgNames) ? getNameFromMapping(ArgNames[I]).c_str() : "unknown");
   }
 #endif
 
   AsyncInfoTy AsyncInfo(Device);
+  AsyncInfoTy &AI = Device.AIM.registerAI(&AsyncInfo);
   int Rc = targetDataBegin(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes,
-                           ArgTypes, ArgNames, ArgMappers, AsyncInfo);
+                           ArgTypes, ArgNames, ArgMappers, AI);
   if (Rc == OFFLOAD_SUCCESS)
-    Rc = AsyncInfo.synchronize();
+    Rc = AI.synchronize(true);
   handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
 }
 
 EXTERN void __tgt_target_data_begin_nowait_mapper(
     ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
     void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
     void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
     void *NoAliasDepList) {
for (int I = 0; I < ArgNum; ++I) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        I, DPxPTR(ArgsBase[I]), DPxPTR(Args[I]), ArgSizes[I], ArgTypes[I],
        (ArgNames) ? getNameFromMapping(ArgNames[I]).c_str() : "unknown");
   }
 #endif
 
   AsyncInfoTy AsyncInfo(Device);
+  AsyncInfoTy &AI = Device.AIM.registerAI(&AsyncInfo);
   int Rc = targetDataEnd(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes,
-                         ArgTypes, ArgNames, ArgMappers, AsyncInfo);
+                         ArgTypes, ArgNames, ArgMappers, AI);
   if (Rc == OFFLOAD_SUCCESS)
-    Rc = AsyncInfo.synchronize();
+    Rc = AI.synchronize(true);
   handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
 }
 
 EXTERN void __tgt_target_data_end_nowait_mapper(
     ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
     void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
     void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
     void *NoAliasDepList) {
EXTERN void __tgt_target_data_update_mapper(ident_t *Loc, int64_t DeviceId,
   }
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(Loc, DeviceId, ArgNum, ArgSizes, ArgTypes, ArgNames,
                          "Updating OpenMP data");
 
   DeviceTy &Device = *PM->Devices[DeviceId];
   AsyncInfoTy AsyncInfo(Device);
+  AsyncInfoTy &AI = Device.AIM.registerAI(&AsyncInfo);
   int Rc = targetDataUpdate(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes,
-                            ArgTypes, ArgNames, ArgMappers, AsyncInfo);
+                            ArgTypes, ArgNames, ArgMappers, AI);
   if (Rc == OFFLOAD_SUCCESS)
-    Rc = AsyncInfo.synchronize();
+    Rc = AI.synchronize(true);
   handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
 }
 
 EXTERN void __tgt_target_data_update_nowait_mapper(
     ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
     void **Args, int64_t *ArgSizes, int64_t *ArgTypes, map_var_info_t *ArgNames,
     void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
     void *NoAliasDepList) {

 #endif
 
   bool IsTeams = NumTeams != -1;
   if (!IsTeams)
     NumTeams = 0;
 
   DeviceTy &Device = *PM->Devices[DeviceId];
   AsyncInfoTy AsyncInfo(Device);
+  AsyncInfoTy &AI = Device.AIM.registerAI(&AsyncInfo);
   int Rc = target(Loc, Device, HostPtr, Args->NumArgs, Args->ArgBasePtrs,
                   Args->ArgPtrs, Args->ArgSizes, Args->ArgTypes, Args->ArgNames,
                   Args->ArgMappers, NumTeams, ThreadLimit, Args->Tripcount,
-                  IsTeams, AsyncInfo);
-  if (Rc == OFFLOAD_SUCCESS)
-    Rc = AsyncInfo.synchronize();
+                  IsTeams, AI);
+  if (Rc == OFFLOAD_SUCCESS) {
+    /// Synchronization is not forced when the env var is enabled and the target
+    /// region does not have any memory transfer
+    Rc = AI.synchronize();
+  }
   handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
   assert(Rc == OFFLOAD_SUCCESS && "__tgt_target_kernel unexpected failure!");
   return OMP_TGT_SUCCESS;
 }
 
 EXTERN int __tgt_target_kernel_nowait(
     ident_t *Loc, int64_t DeviceId, int32_t NumTeams, int32_t ThreadLimit,
     void *HostPtr, __tgt_kernel_arguments *Args, int32_t DepNum, void *DepList,

openmp/libomptarget/src/omptarget.cpp

 
 #include "omptarget.h"
 #include "device.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
 #include <cstdint>
+#include <utility>
 #include <vector>
 
 using llvm::SmallVector;
 
-int AsyncInfoTy::synchronize() {
+bool AsyncFlag = false;
+
+int AsyncInfoTy::synchronize(bool ForceSync) {
   int Result = OFFLOAD_SUCCESS;
+  if (!AsyncFlag || ForceSync) {
     if (AsyncInfo.Queue) {
+      DP("Device %d synchronization\n", Device.DeviceID);
       // If we have a queue we need to synchronize it now.
       Result = Device.synchronize(*this);
       assert(AsyncInfo.Queue == nullptr &&
              "The device plugin should have nulled the queue to indicate there "
              "are no outstanding actions!");
     }
+  }
   return Result;
 }
 
 void *&AsyncInfoTy::getVoidPtrLocation() {
   BufferLocations.push_back(nullptr);
   return BufferLocations.back();
 }
 
+/// Async info manager
+AsyncInfoMng::AsyncInfoMng(DeviceTy &Device) : Device(Device) {}
+
+AsyncInfoMng::~AsyncInfoMng() {
+  for (const auto &It : AsyncInfoM)
+    delete (It.second);
+  AsyncInfoM.clear();
+}
+
+AsyncInfoTy &AsyncInfoMng::registerAI(AsyncInfoTy *AI) {
+  if (!AsyncFlag)
+    return *AI;
+
+  /// Get async info
+  std::lock_guard<std::mutex> MapLock(AsyncMtx);
+  llvm::DenseMap<std::thread::id, AsyncInfoTy *>::iterator it =
+      AsyncInfoM.find(std::this_thread::get_id());
+  if (it != AsyncInfoM.end())
+    return *(it->second);
+  /// Add new Async Info
+  AsyncInfoTy *AsyncInfo = new AsyncInfoTy(Device, false);
+  AsyncInfoM.insert(std::make_pair(std::this_thread::get_id(), AsyncInfo));
+  return *AsyncInfo;
+}
+
 /* All begin addresses for partially mapped structs must be 8-aligned in order
  * to ensure proper alignment of members. E.g.
  *
  * struct S {
  *   int a;   // 4-aligned
  *   int b;   // 4-aligned
  *   int *p;  // 8-aligned
  * } s1;
  * ...
  * #pragma omp target map(tofrom: s1.b, s1.p[0:N])
  * {
  *   s1.b = 5;
  *   for (int i...) s1.p[i] = ...;
  * }
  *
  * Here we are mapping s1 starting from member b, so BaseAddress=&s1=&s1.a and
  * BeginAddress=&s1.b. Let's assume that the struct begins at address 0x100,
  * then &s1.a=0x100, &s1.b=0x104, &s1.p=0x108. Each member obeys the alignment
  * requirements for its type. Now, when we allocate memory on the device, in
  * CUDA's case cuMemAlloc() returns an address which is at least 256-aligned.
  * This means that the chunk of the struct on the device will start at a
  * 256-aligned address, let's say 0x200. Then the address of b will be 0x200 and
  * address of p will be a misaligned 0x204 (on the host there was no need to add
for (auto &Lib : Device.PendingCtorsDtors) {
           }
         }
         // Clear the list to indicate that this device has been used
         Lib.second.PendingCtors.clear();
         DP("Done with pending ctors for lib " DPxMOD "\n", DPxPTR(Lib.first));
       }
     }
     // All constructors have been issued, wait for them now.
-    if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
+    if (AsyncInfo.synchronize(true) != OFFLOAD_SUCCESS)
       return OFFLOAD_FAIL;
   }
   Device.HasPendingGlobals = false;
 
   return OFFLOAD_SUCCESS;
 }
 
 void handleTargetOutcome(bool Success, ident_t *Loc) {

 
 } // namespace
 
 /// Internal function to undo the mapping and retrieve the data from the device.
 int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, map_var_info_t *ArgNames,
                   void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) {
-  int Ret;
+  int Ret = OFFLOAD_SUCCESS;
   SmallVector<PostProcessingInfo> PostProcessingPtrs;
   void *FromMapperBase = nullptr;
   // process each input.
   for (int32_t I = ArgNum - 1; I >= 0; --I) {
     // Ignore private variables and arrays - there is no mapping for them.
     // Also, ignore the use_device_ptr directive, it has no effect here.
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
         (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       }
 
       // Add pointer to the buffer for post-synchronize processing.
       PostProcessingPtrs.emplace_back(HstPtrBegin, DataSize, ArgTypes[I],
                                       DelEntry && !IsHostPtr, TPR);
     }
   }
 
-  // TODO: We should not synchronize here but pass the AsyncInfo object to the
-  //       allocate/deallocate device APIs.
-  //
-  // We need to synchronize before deallocating data.
-  Ret = AsyncInfo.synchronize();
-  if (Ret != OFFLOAD_SUCCESS)
-    return OFFLOAD_FAIL;
-
   // Deallocate target pointer
   for (PostProcessingInfo &Info : PostProcessingPtrs) {
     // If we marked the entry to be deleted we need to verify no other thread
     // reused it by now. If deletion is still supposed to happen by this thread
     // LR will be set and exclusive access to the HDTT map will avoid another
     // thread reusing the entry now. Note that we do not request (exclusive)
     // access to the HDTT map if Info.DelEntry is not set.
     LookupResult LR;
     DeviceTy::HDTTMapAccessorTy HDTTMap =
         Device.HostDataToTargetMap.getExclusiveAccessor(!Info.DelEntry);
-
     if (Info.DelEntry) {
       LR = Device.lookupMapping(HDTTMap, Info.HstPtrBegin, Info.DataSize);
       if (LR.Entry->getTotalRefCount() != 0 ||
           LR.Entry->getDeleteThreadId() != std::this_thread::get_id()) {
         // The thread is not in charge of deletion anymore. Give up access to
         // the HDTT map and unset the deletion flag.
         HDTTMap.destroy();
         Info.DelEntry = false;
auto CB = [&](ShadowPtrListTy::iterator &Itr) {
         Device.ShadowPtrMap.erase(OldItr);
       } else {
         ++Itr;
       }
       return OFFLOAD_SUCCESS;
     };
     applyToShadowMapEntries(Device, CB, Info.HstPtrBegin, Info.DataSize,
                             Info.TPR);
-
     // If we are deleting the entry the DataMapMtx is locked and we own the
     // entry.
     if (Info.DelEntry) {
       if (!FromMapperBase || FromMapperBase != Info.HstPtrBegin)
-        Ret = Device.deallocTgtPtr(HDTTMap, LR, Info.DataSize);
+        Ret = Device.deallocTgtPtr(HDTTMap, LR, Info.DataSize, AsyncInfo);
 
       if (Ret != OFFLOAD_SUCCESS) {
         REPORT("Deallocating data from device failed.\n");
         break;
       }
     }
   }
 
if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Copying data from device failed.\n");
       return OFFLOAD_FAIL;
     }
 
     auto CB = [&](ShadowPtrListTy::iterator &Itr) {
       void **ShadowHstPtrAddr = (void **)Itr->first;
       // Wait for device-to-host memcopies for whole struct to complete,
       // before restoring the correct host pointer.
-      if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
+      if (AsyncInfo.synchronize(true) != OFFLOAD_SUCCESS)
         return OFFLOAD_FAIL;
       *ShadowHstPtrAddr = Itr->second.HstPtrVal;
       DP("Restoring original host pointer value " DPxMOD
          " for host pointer " DPxMOD "\n",
          DPxPTR(Itr->second.HstPtrVal), DPxPTR(ShadowHstPtrAddr));
       ++Itr;
       return OFFLOAD_SUCCESS;
     };

openmp/libomptarget/src/rtl.cpp

static const char *RTLNames[] = {
     /* SX-Aurora VE target  */ "libomptarget.rtl.ve.so",
     /* AMDGPU target        */ "libomptarget.rtl.amdgpu.so",
     /* Remote target        */ "libomptarget.rtl.rpc.so",
 };
 
 PluginManager *PM;
 
 static char *ProfileTraceFile = nullptr;
+extern bool AsyncFlag;
 
 __attribute__((constructor(101))) void init() {
   DP("Init target library!\n");
 
   bool UseEventsForAtomicTransfers = true;
   if (const char *ForceAtomicMap = getenv("LIBOMPTARGET_MAP_FORCE_ATOMIC")) {
     std::string ForceAtomicMapStr(ForceAtomicMap);
     if (ForceAtomicMapStr == "false" || ForceAtomicMapStr == "FALSE")
       UseEventsForAtomicTransfers = false;
     else if (ForceAtomicMapStr != "true" && ForceAtomicMapStr != "TRUE")
       fprintf(stderr,
               "Warning: 'LIBOMPTARGET_MAP_FORCE_ATOMIC' accepts only "
               "'true'/'TRUE' or 'false'/'FALSE' as options, '%s' ignored\n",
               ForceAtomicMap);
   }
 
   PM = new PluginManager(UseEventsForAtomicTransfers);
 
   ProfileTraceFile = getenv("LIBOMPTARGET_PROFILE");
   // TODO: add a configuration option for time granularity
   if (ProfileTraceFile)
     timeTraceProfilerInitialize(500 /* us */, "libomptarget");
+
+  /// Asynchronous execution flag
+  if (char *EnvStr = getenv("LIBOMPTARGET_INTRA_THREAD_ASYNC"))
+    AsyncFlag = std::stoi(EnvStr) ? true : false;
+  DP("Asynchronous execution %s\n", AsyncFlag ? "Enabled" : "Disabled");
 }
 
 __attribute__((destructor(101))) void deinit() {
   DP("Deinit target library!\n");
 
   for (auto *R : PM->RTLs.UsedRTLs) {
     // Plugins can either destroy their local state using global variables
     // or attribute(destructor) functions or by implementing deinit_plugin
for (auto *R : UsedRTLs) {
                             AsyncInfo);
             if (Rc != OFFLOAD_SUCCESS) {
               DP("Running destructor " DPxMOD " failed.\n", DPxPTR(Dtor));
             }
           }
           // Remove this library's entry from PendingCtorsDtors
           Device.PendingCtorsDtors.erase(Desc);
           // All constructors have been issued, wait for them now.
-          if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
+          if (AsyncInfo.synchronize(true) != OFFLOAD_SUCCESS)
             DP("Failed synchronizing destructors kernels.\n");
         }
         Device.PendingGlobalsMtx.unlock();
       }
 
       DP("Unregistered image " DPxMOD " from RTL " DPxMOD "!\n",
          DPxPTR(Img->ImageStart), DPxPTR(R->LibraryHandler.get()));