OpenMP asynchronous memory copy support

[jz10]

Summary

We introduced the implementation of supporting asynchronous routines with depend objects specified in Version 5.1 of the OpenMP Application Programming Interface. In brief, these routines omp_target_memcpy_async and omp_target_memcpy_rect_async perform asynchronous (nonblocking) memory copies between any
combination of host and device pointers. The basic idea is to create the implicit tasks to carry the memory copy calls and handle dependencies specified by depend objects. The implicit tasks are executed via hidden helper thread in OpenMP runtime.

[josemonsalve2]

Thanks for implementing this. I have added a comment inlined.

openmp/libomptarget/src/private.h

119

This would be the third location where this struct is duplicated: interop.h, kmp.h and this file. Would it make sense to try to add it to another common header file?

[jdoerfert]

First set of comments.

clang/docs/ReleaseNotes.rst

253 ↗

(On Diff #468290)

Unrelated, probably bad diff origin.

openmp/libomptarget/src/api.cpp

213

static also below.

217

LLVM codeing style, please. So here, Args

220

not 0 but nullptr, also lots of other places.

226

Why ignore the return value of memcpy and return 0 here?

242

5? Why 5?

247

Proper names, please. Also coding style.

252

This is very suspicious. Why can't we have a kmp_tasking_flags_t object?

265

Does this return anything? The use of Rc seems useless. Why do you access args_ for some parts and not for others? That said, where does the hidden helper need access to the dependences anyway?

436

This screams helper as it is literally the same code modulo 5 characters in a few places. Please refactor and run clang-format on the patch.

[jz10]

Thanks Johannes for your comments, and I relied them below

1. format issues

I ran clang-format to reformat, please check if there's any missed things;

2. replace '0' with 'nullptr'

fixed

3. proper return value for helper functions and async functions

fixed

4. Why can't we have a kmp_tasking_flags_t object?

I followed the same access approach in openmp/runtime , so I didn't change this part. But I can revise if it is needed

5. make helper function as 'static'

fixed

6. Why do you access args_ for some parts and not for others? That said, where does the hidden helper need access to the dependences anyway?

there's type cast for depend objects from 'omp_depend_t' to 'kmp_depend_info_t*', and the array of casted depend objects is consumed by '__kmpc_omp_task_with_deps' , to make it safe, I just make larray of casted depend objects to live longer, thus attached it to Args object.

[jdoerfert]

6. Why do you access args_ for some parts and not for others? That said, where does the hidden helper need access to the dependences anyway?

there's type cast for depend objects from 'omp_depend_t' to 'kmp_depend_info_t*', and the array of casted depend objects is consumed by '__kmpc_omp_task_with_deps' , to make it safe, I just make larray of casted depend objects to live longer, thus attached it to Args object.

So, you are saying the task_with_deps function does *not* copy the dependences and therefore the array has to outlive the function?

openmp/libomptarget/src/private.h

258

As mentioned before. There is a lot of literal duplication here with no obvious benefit. Use a base class, or other schemes, to cut down on that. Also, the functions further above should share the same code.

[jz10]

I'm not sure if it copies, will check it to confirm

[tianshilei1992]

openmp/libomptarget/src/api.cpp

262

Why does it need a _ at the end?

407

I'm not sure if it is a failure if source and destination are same.

openmp/libomptarget/src/private.h

119

IIRC, there are some (non-technical) issues on using kmp.h in libomptarget.

181

proxy task is not needed in this patch. we don't have detached task.

222

This is really not LLVM code style.

[tianshilei1992]

openmp/libomptarget/src/api.cpp

217

I didn't see Args is freed after use. There is memory leak.

openmp/libomptarget/src/private.h

197

argument name is not in LLVM style

[jz10]

Thanks Johannes and Shilei

1. "So, you are saying the task_with_deps function does *not* copy the dependences and therefore the array has to outlive the function?"

I checked the omp_task_with_deps, it does some copy operations , so I move the depobj_list related code out of the TargetMemcpyArgsTy and TargetMemcpyRectArgsTy objects now. Those two classes are also redefined as 'struct'.

2. "> This would be the third location where this struct is duplicated: interop.h, kmp.h and this file. Would it make sense to try to add it to another common header file?

IIRC, there are some (non-technical) issues on using kmp.h in libomptarget."
So should I keep the data structure definitions in private.h? or create a new header file?

3. no release of Args

fixed

4. "I'm not sure if it is a failure if source and destination are same"

we thought about this, and should allow user to do this

5. "This is really not LLVM code style."

I changed the 'class' definition to 'struct' and reformat them. Please check if that works

6. few formatting issues

fixed

[jz10]

one minors fix, i.e. delete the depobj_list from TargetMemcpyArgsTy struct

[jdoerfert]

openmp/libomptarget/src/api.cpp

273

Use an llvm::SmallVector

436

^^^

[jz10]

Thanks Johannes

1. use SmallVector

fixed

2. "module 5 characters"

ran clang-format , please check it that works

[jz10]

1. use SmallVector

fixed

2. "5 characters"

ran clang-format, please check if that works

[jz10]

redo the proper patch, i.e. get rid of content related to clang/doc

[jdoerfert]

Thanks for the latest updates. The duplication is still my main concern here. See below.

openmp/libomptarget/src/api.cpp

373

My point about the "5 character difference" was that this function is basically the same as __kmpc_target_memcpy_async_helper and similarly omp_target_memcpy_rect_async is almost the same as omp_target_memcpy_async. Can we use common helper functions to avoid duplicating all that code twice?
(Put differently: I assume that you implemented one of these functions and then copied it and renamed a few things to get the other. Doing that should have made it obvious that there is duplication and we want to avoid duplication.)

439

That's not how vectors work. push_back(...)

openmp/libomptarget/src/private.h

228

These structs need doxygen comments describing where they are used.

[jz10]

Thanks Johannes and Shilei

1. using a common helper function

fixed

2. using push_back for SmallVector

fixed

3. add doxygen comments for struct

added, please check if that works

[jdoerfert]

openmp/libomptarget/src/api.cpp

425

Now we have a first helper for the "async" part, we still should deduplicate the entry point code. Above, all but lines 385-387 are the same as in omp_target_memcpy_async. Can we also not duplicate those lines?

In this code there are also various places with variables not named according to the style guide, e.g. _Count, and with names not helpful, e.g. ErrXYZ. Please address those too. Similarly, use nullptr instead of NULL.

The problem with the int32 Flags I mentioned already. Let's avoid type punning and create the object with the right type from the very beginning.

[tianshilei1992]

openmp/libomptarget/src/private.h

107

Can we put all KMP related code into a separate header, but of course not called kmp.h? I don't think it's a good idea to have them everywhere. That's quite unfortunate that we have to duplicate the code.

229

Having a variable suffix with _ is generally not a good coding style, not LLVM style. You can potentially use the same name. The compiler can tell them apart. For those values, such as NumDims, you could potentially do int NumDims = 0 when defining them.

[jz10]

Thanks Johannes and Shilei

1. '385-387 are the same as in omp_target_memcpy_async. Can we also not duplicate those lines?'

I put the common code (i.e.helper task creation) into another static function

2. 'In this code there are also various places with variables not named according to the style guide'

fixed, please check if there's remaining issues

3. 'The problem with the int32 Flags I mentioned already. '

The flag variable was defined as 'kmp_int32', since its consumer '__kmpc_omp_target_task_alloc' needs 'kmp_int32' type as input. the type cast to 'kmp_tasking_flags_t' is to set the 'hidden_helper' bit. So it seems that there's no better option for us. Please let me know your suggestion.

4. "Can we put all KMP related code into a separate header"

we used both kmp relevent data structure/types and APIs, so should I wrap all those relevant code into several tool functions and put them into separate header file?

5. 'aving a variable suffix with _ is generally not a good coding style'

fixed

[jdoerfert]

In D136103#3876097, @jz10 wrote:

Thanks Johannes and Shilei

1. '385-387 are the same as in omp_target_memcpy_async. Can we also not duplicate those lines?'

I put the common code (i.e.helper task creation) into another static function

I feel like I'm missing something. As said before, all but 3 lines are identical in these two functions.
Now you created a helper for 1/3 of those identical lines but left the other 2/3 being duplicated. Could
you elaborate why?

2. 'In this code there are also various places with variables not named according to the style guide'

fixed, please check if there's remaining issues

3. 'The problem with the int32 Flags I mentioned already. '

The flag variable was defined as 'kmp_int32', since its consumer '__kmpc_omp_target_task_alloc' needs 'kmp_int32' type as input. the type cast to 'kmp_tasking_flags_t' is to set the 'hidden_helper' bit. So it seems that there's no better option for us. Please let me know your suggestion.

Yeah,.. the API is broken. Let's keep it as is. Apologies to have brought it up.

4. "Can we put all KMP related code into a separate header"

we used both kmp relevent data structure/types and APIs, so should I wrap all those relevant code into several tool functions and put them into separate header file?

5. 'aving a variable suffix with _ is generally not a good coding style'

fixed

[jz10]

Thanks Johannes

1. 'I feel like I'm missing something. As said before, all but 3 lines are identical in these two functions.

Now you created a helper for 1/3 of those identical lines but left the other 2/3 being duplicated. Could
you elaborate why?'
I put the common part into task_creation function

[jdoerfert]

Looks good from my end. Thanks for the updates, I hope they make sense in the end.

Some minor style comments below. @tianshilei1992 Should accept once the header concerns are resolved.

openmp/libomptarget/src/api.cpp

263–264

You can delete a nullptr.

274–279

281–283

[jz10]

Thanks Johannes
I revised those issues, please check if those work

[tianshilei1992]

we used both kmp relevent data structure/types and APIs, so should I wrap all those relevant code into several tool functions and put them into separate header file?

IMO we can put all KMP related code into one header and include it where needed. For the others, the API header has to be included anyway.

[jz10]

Regarding this, can we just move those two helper functions to private.h ?

[jz10]

I checked through private.h, this header actually does the functionality
that contains all kmp and kmpc related data structures and APIs, so should
we still have to split a separate header file?

[tianshilei1992]

In D136103#3880651, @jz10 wrote:

I checked through private.h, this header actually does the functionality
that contains all kmp and kmpc related data structures and APIs, so should
we still have to split a separate header file?

Yeah you are right. Nah, that's fine. Let it be for now.

Can you add a couple of tests?

openmp/libomptarget/src/api.cpp

250

potentially an auto *Fn is more appropriate here.

[jz10]

Sure, where should I add those tests?

[tianshilei1992]

openmp/libomptarget/test/api/ is where we usually tests those APIs.

[jz10]

Thanks Johannes and Shilei

I added few test cases for asynchronous routine at test/api folder

[tianshilei1992]

openmp/libomptarget/test/api/omp_target_memcpy_async1.c

2

There is no RUN line here so the test will not be triggered. Can you refer to other test cases to see how a test is run?

[jz10]

Thanks Shilei

Add RUN line for each test cases

[tianshilei1992]

openmp/libomptarget/test/api/omp_target_memcpy_async1.c

3

Does it work on AMDGPU and other targets? Why does it require Nvidia here?

openmp/libomptarget/test/api/omp_target_memcpy_rect_async3.c

21 ↗

(On Diff #470354)

If I read this test correctly, this test is more for performance evaluation instead of correctness. We don't have any performance test cases yet, and I'm not sure we need them right now.

181 ↗

(On Diff #470354)

If we don't profile the time, I don't see a point in keeping them.

[jz10]

Thanks Shilei

1. "Does it work on AMDGPU and other targets? Why does it require Nvidia here?"

No, I just remove this line

2. "We don't have any performance test cases yet, and I'm not sure we need them right now."

Yes, I remove those two performance test cases

[jz10]

hi Johannes and Shilei, is there revision that needs to be done on this
patch? please let me know

[tianshilei1992]

Generally looks good to me. Can you check all resolved comments to make sure there is no open comments?

openmp/libomptarget/src/api.cpp

213

same here

245

Since this function is not part of libomp and it's not gonna be an interface function, no need to name it as __kmpc.

[jz10]

Thanks Shilei

1. "Can you check all resolved comments to make sure there is no open comments?"

checked through the comments you and Johannes made, no more issues

2. "Since this function is not part of libomp and it's not gonna be an interface function, no need to name it as __kmpc."

revised helper functions' names

[jdoerfert]

LG, thanks for all the changes.

[tianshilei1992]

Thanks! LG.

[jdoerfert]

Any reason this hasn't been pushed yet?

[jplehr]

Rebase to make ready for land

[jplehr]

Rebased and enabled tests for generic devices.
Resulted in one test failure

Failed Tests (3):
libomptarget :: amdgcn-amd-amdhsa :: api/omp_target_memcpy_rect_async1.c
libomptarget :: x86_64-pc-linux-gnu :: api/omp_target_memcpy_rect_async1.c
libomptarget :: x86_64-pc-linux-gnu-LTO :: api/omp_target_memcpy_rect_async1.c#

[jplehr]

Fix bug to corectly support the maximally supported dimensions as required by the spec.

[jplehr]

Removed accidentally added code

[jdoerfert]

LG

openmp/libomptarget/src/api.cpp

386

!DST && !SRC

openmp/libomptarget/include/interop.h

 omp_get_interop_type_desc(const omp_interop_t, omp_interop_property_t);
 /*!
  * The `omp_get_interop_rc_desc` routine retrieves a description of the return
  * code associated with an `omp_interop_t` object.
  */
 extern const char *__KAI_KMPC_CONVENTION
 omp_get_interop_rc_desc(const omp_interop_t, omp_interop_rc_t);
 
-typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */
-  /* Compiler flags */             /* Total compiler flags must be 16 bits */
-  unsigned tiedness : 1;           /* task is either tied (1) or untied (0) */
-  unsigned final : 1;              /* task is final(1) so execute immediately */
-  unsigned merged_if0 : 1; // no __kmpc_task_{begin/complete}_if0 calls in if0
-  unsigned destructors_thunk : 1; // set if the compiler creates a thunk to
-  unsigned proxy : 1; // task is a proxy task (it will be executed outside the
-  unsigned priority_specified : 1; // set if the compiler provides priority
-  unsigned detachable : 1;         // 1 == can detach */
-  unsigned unshackled : 1;         /* 1 == unshackled task */
-  unsigned target : 1;             /* 1 == target task */
-  unsigned reserved : 7;           /* reserved for compiler use */
-  unsigned tasktype : 1;    /* task is either explicit(1) or implicit (0) */
-  unsigned task_serial : 1; // task is executed immediately (1) or deferred (0)
-  unsigned tasking_ser : 1; // all tasks in team are either executed immediately
-  unsigned team_serial : 1; // entire team is serial (1) [1 thread] or parallel
-  unsigned started : 1;     /* 1==started, 0==not started     */
-  unsigned executing : 1;   /* 1==executing, 0==not executing */
-  unsigned complete : 1;    /* 1==complete, 0==not complete   */
-  unsigned freed : 1;       /* 1==freed, 0==allocated        */
-  unsigned native : 1;      /* 1==gcc-compiled task, 0==intel */
-  unsigned reserved31 : 7;  /* reserved for library use */
-} kmp_tasking_flags_t;
-
 typedef enum omp_interop_backend_type_t {
   // reserve 0
   omp_interop_backend_type_cuda_1 = 1,
 } omp_interop_backend_type_t;
 
 typedef enum kmp_interop_type_t {
   kmp_interop_type_unknown = -1,
   kmp_interop_type_platform,

openmp/libomptarget/src/api.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "device.h"
 #include "omptarget.h"
 #include "private.h"
 #include "rtl.h"
 
+#include "llvm/ADT/SmallVector.h"
+
 #include <climits>
 #include <cstdlib>
 #include <cstring>
 #include <mutex>
 
 EXTERN int omp_get_num_devices(void) {
   TIMESCOPE();
   PM->RTLsMtx.lock();
if (SrcDevice == omp_get_initial_device() &&
     }
     free(Buffer);
   }
 
   DP("omp_target_memcpy returns %d\n", Rc);
   return Rc;
 }
 
+// The helper function that calls omp_target_memcpy or omp_target_memcpy_rect
+static int libomp_target_memcpy_async_helper(kmp_int32 Gtid, kmp_task_t *Task) {
+  if (Task == nullptr)
+    return OFFLOAD_FAIL;
+
+  TargetMemcpyArgsTy *Args = (TargetMemcpyArgsTy *)Task->shareds;
+
+  if (Args == nullptr)
+    return OFFLOAD_FAIL;
+
+  // Call blocked version
+  int Rc = OFFLOAD_SUCCESS;
+  if (Args->IsRectMemcpy) {
+    Rc = omp_target_memcpy_rect(
+        Args->Dst, Args->Src, Args->ElementSize, Args->NumDims, Args->Volume,
+        Args->DstOffsets, Args->SrcOffsets, Args->DstDimensions,
+        Args->SrcDimensions, Args->DstDevice, Args->SrcDevice);
+
+    DP("omp_target_memcpy_rect returns %d\n", Rc);
+  } else {
+    Rc = omp_target_memcpy(Args->Dst, Args->Src, Args->Length, Args->DstOffset,
+                           Args->SrcOffset, Args->DstDevice, Args->SrcDevice);
+
+    DP("omp_target_memcpy returns %d\n", Rc);
+  }
+
+  // Release the arguments object
+  delete Args;
+
+  return Rc;
+}
+
+// Allocate and launch helper task
+static int libomp_helper_task_creation(TargetMemcpyArgsTy *Args,
+                                       int DepObjCount,
+                                       omp_depend_t *DepObjList) {
+  // Create global thread ID
+  int Gtid = __kmpc_global_thread_num(nullptr);
+  int (*Fn)(kmp_int32, kmp_task_t *) = &libomp_target_memcpy_async_helper;
+
+  // Setup the hidden helper flags;
+  kmp_int32 Flags = 0;
+  kmp_tasking_flags_t *InputFlags = (kmp_tasking_flags_t *)&Flags;
+  InputFlags->hidden_helper = 1;
+
+  // Alloc helper task
+  kmp_task_t *Ptr = __kmpc_omp_target_task_alloc(nullptr, Gtid, Flags,
+                                                 sizeof(kmp_task_t), 0, Fn, -1);
+
+  if (Ptr == nullptr) {
+    // Task allocation failed, delete the argument object
+    delete Args;
+
+    return OFFLOAD_FAIL;
+  }
+
+  // Setup the arguments passed to helper task
+  Ptr->shareds = Args;
+
+  // Convert the type of depend objects
+  llvm::SmallVector<kmp_depend_info_t> DepObjs;
+  for (int i = 0; i < DepObjCount; i++) {
+    omp_depend_t DepObj = DepObjList[i];
+    DepObjs.push_back(*((kmp_depend_info_t *)DepObj));
+  }
+
+  // Launch the helper task
+  int Rc = __kmpc_omp_task_with_deps(nullptr, Gtid, Ptr, DepObjCount,
+                                     DepObjs.data(), 0, nullptr);
+
+  return Rc;
+}
+
+EXTERN int omp_target_memcpy_async(void *Dst, const void *Src, size_t Length,
+                                   size_t DstOffset, size_t SrcOffset,
+                                   int DstDevice, int SrcDevice,
+                                   int DepObjCount, omp_depend_t *DepObjList) {
+  TIMESCOPE();
+  DP("Call to omp_target_memcpy_async, dst device %d, src device %d, "
+     "dst addr " DPxMOD ", src addr " DPxMOD ", dst offset %zu, "
+     "src offset %zu, length %zu\n",
+     DstDevice, SrcDevice, DPxPTR(Dst), DPxPTR(Src), DstOffset, SrcOffset,
+     Length);
+
+  // Check the source and dest address
+  if (Dst == nullptr || Src == nullptr)
+    return OFFLOAD_FAIL;
+
+  // Create task object
+  TargetMemcpyArgsTy *Args = new TargetMemcpyArgsTy(
+      Dst, Src, Length, DstOffset, SrcOffset, DstDevice, SrcDevice);
+
+  // Create and launch helper task
+  int Rc = libomp_helper_task_creation(Args, DepObjCount, DepObjList);
+
+  DP("omp_target_memcpy_async returns %d\n", Rc);
+  return Rc;
+}
+
 EXTERN int
 omp_target_memcpy_rect(void *Dst, const void *Src, size_t ElementSize,
                        int NumDims, const size_t *Volume,
                        const size_t *DstOffsets, const size_t *SrcOffsets,
                        const size_t *DstDimensions, const size_t *SrcDimensions,
                        int DstDevice, int SrcDevice) {
   TIMESCOPE();
   DP("Call to omp_target_memcpy_rect, dst device %d, src device %d, "
for (size_t I = 0; I < Volume[0]; ++I) {
       }
     }
   }
 
   DP("omp_target_memcpy_rect returns %d\n", Rc);
   return Rc;
 }
 
+EXTERN int omp_target_memcpy_rect_async(
+    void *Dst, const void *Src, size_t ElementSize, int NumDims,
+    const size_t *Volume, const size_t *DstOffsets, const size_t *SrcOffsets,
+    const size_t *DstDimensions, const size_t *SrcDimensions, int DstDevice,
+    int SrcDevice, int DepObjCount, omp_depend_t *DepObjList) {
+  TIMESCOPE();
+  DP("Call to omp_target_memcpy_rect_async, dst device %d, src device %d, "
+     "dst addr " DPxMOD ", src addr " DPxMOD ", dst offsets " DPxMOD ", "
+     "src offsets " DPxMOD ", dst dims " DPxMOD ", src dims " DPxMOD ", "
+     "volume " DPxMOD ", element size %zu, num_dims %d\n",
+     DstDevice, SrcDevice, DPxPTR(Dst), DPxPTR(Src), DPxPTR(DstOffsets),
+     DPxPTR(SrcOffsets), DPxPTR(DstDimensions), DPxPTR(SrcDimensions),
+     DPxPTR(Volume), ElementSize, NumDims);
+
+  // Need to check this first to not return OFFLOAD_FAIL instead
+  if (!Dst && !Src) {
+    DP("Call to omp_target_memcpy_rect returns max supported dimensions %d\n",
+       INT_MAX);
+    return INT_MAX;
+  }
+
+  // Check the source and dest address
+  if (Dst == nullptr || Src == nullptr)
+    return OFFLOAD_FAIL;
+
+  // Create task object
+  TargetMemcpyArgsTy *Args = new TargetMemcpyArgsTy(
+      Dst, Src, ElementSize, NumDims, Volume, DstOffsets, SrcOffsets,
+      DstDimensions, SrcDimensions, DstDevice, SrcDevice);
+
+  // Create and launch helper task
+  int Rc = libomp_helper_task_creation(Args, DepObjCount, DepObjList);
+
+  DP("omp_target_memcpy_rect_async returns %d\n", Rc);
+  return Rc;
+}
+
 EXTERN int omp_target_associate_ptr(const void *HostPtr, const void *DevicePtr,
                                     size_t Size, size_t DeviceOffset,
                                     int DeviceNum) {
   TIMESCOPE();
   DP("Call to omp_target_associate_ptr with host_ptr " DPxMOD ", "
      "device_ptr " DPxMOD ", size %zu, device_offset %zu, device_num %d\n",
      DPxPTR(HostPtr), DPxPTR(DevicePtr), Size, DeviceOffset, DeviceNum);
 
   if (!HostPtr || !DevicePtr || Size <= 0) {
     REPORT("Call to omp_target_associate_ptr with invalid arguments\n");
     return OFFLOAD_FAIL;
   }
 
   if (DeviceNum == omp_get_initial_device()) {
     REPORT("omp_target_associate_ptr: no association possible on the host\n");
     return OFFLOAD_FAIL;
   }
 
   if (!deviceIsReady(DeviceNum)) {
     REPORT("omp_target_associate_ptr returns OFFLOAD_FAIL\n");
     return OFFLOAD_FAIL;
   }
 
   DeviceTy &Device = *PM->Devices[DeviceNum];
   void *DeviceAddr = (void *)((uint64_t)DevicePtr + (uint64_t)DeviceOffset);
   int Rc = Device.associatePtr(const_cast<void *>(HostPtr),
                                const_cast<void *>(DeviceAddr), Size);
   DP("omp_target_associate_ptr returns %d\n", Rc);
   return Rc;
 }
 
 EXTERN int omp_target_disassociate_ptr(const void *HostPtr, int DeviceNum) {
   TIMESCOPE();
   DP("Call to omp_target_disassociate_ptr with host_ptr " DPxMOD ", "
      "device_num %d\n",
      DPxPTR(HostPtr), DeviceNum);
 
   if (!HostPtr) {
     REPORT("Call to omp_target_associate_ptr with invalid host_ptr\n");
     return OFFLOAD_FAIL;

openmp/libomptarget/src/exports

global:
     omp_get_num_devices;
     omp_get_device_num;
     omp_get_initial_device;
     omp_target_alloc;
     omp_target_free;
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;
+    omp_target_memcpy_async;
+    omp_target_memcpy_rect_async;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
     llvm_omp_target_free_host;
     llvm_omp_target_free_shared;
     llvm_omp_target_free_device;

openmp/libomptarget/src/private.h

 extern "C" {
 #endif
 /*!
  * The ident structure that describes a source location.
  * The struct is identical to the one in the kmp.h file.
  * We maintain the same data structure for compatibility.
  */
 typedef int kmp_int32;
+typedef int64_t kmp_int64;
 typedef intptr_t kmp_intptr_t;
+
+typedef void *omp_depend_t;
+struct kmp_task;
+typedef kmp_int32 (*kmp_routine_entry_t)(kmp_int32, struct kmp_task *);
+typedef struct kmp_task {
+  void *shareds;
+  kmp_routine_entry_t routine;
+  kmp_int32 part_id;
+} kmp_task_t;
+
+typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */
+  /* Compiler flags */             /* Total compiler flags must be 16 bits */
+  unsigned tiedness : 1;           /* task is either tied (1) or untied (0) */
+  unsigned final : 1;              /* task is final(1) so execute immediately */
+  unsigned merged_if0 : 1; /* no __kmpc_task_{begin/complete}_if0 calls in if0
+                              code path */
+  unsigned destructors_thunk : 1; /* set if the compiler creates a thunk to
+                                     invoke destructors from the runtime */
+  unsigned proxy : 1; /* task is a proxy task (it will be executed outside the
+                         context of the RTL) */
+  unsigned priority_specified : 1; /* set if the compiler provides priority
+                                      setting for the task */
+  unsigned detachable : 1;         /* 1 == can detach */
+  unsigned hidden_helper : 1;      /* 1 == hidden helper task */
+  unsigned reserved : 8;           /* reserved for compiler use */
+
+  /* Library flags */       /* Total library flags must be 16 bits */
+  unsigned tasktype : 1;    /* task is either explicit(1) or implicit (0) */
+  unsigned task_serial : 1; // task is executed immediately (1) or deferred (0)
+  unsigned tasking_ser : 1; // all tasks in team are either executed immediately
+  // (1) or may be deferred (0)
+  unsigned team_serial : 1; // entire team is serial (1) [1 thread] or parallel
+  // (0) [>= 2 threads]
+  /* If either team_serial or tasking_ser is set, task team may be NULL */
+  /* Task State Flags: */
+  unsigned started : 1;    /* 1==started, 0==not started     */
+  unsigned executing : 1;  /* 1==executing, 0==not executing */
+  unsigned complete : 1;   /* 1==complete, 0==not complete   */
+  unsigned freed : 1;      /* 1==freed, 0==allocated        */
+  unsigned native : 1;     /* 1==gcc-compiled task, 0==intel */
+  unsigned reserved31 : 7; /* reserved for library use */
+} kmp_tasking_flags_t;
+
 // Compiler sends us this info:
 typedef struct kmp_depend_info {
   kmp_intptr_t base_addr;
   size_t len;
   struct {
     bool in : 1;
     bool out : 1;
     bool mtx : 1;
   } flags;
 } kmp_depend_info_t;
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_global_thread_num(void *) __attribute__((weak));
 int __kmpc_get_target_offload(void) __attribute__((weak));
 void __kmpc_omp_wait_deps(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 ndeps,
                           kmp_depend_info_t *dep_list, kmp_int32 ndeps_noalias,
                           kmp_depend_info_t *noalias_dep_list)
     __attribute__((weak));
 void **__kmpc_omp_get_target_async_handle_ptr(kmp_int32 gtid)
     __attribute__((weak));
 bool __kmpc_omp_has_task_team(kmp_int32 gtid) __attribute__((weak));
+kmp_task_t *__kmpc_omp_task_alloc(ident_t *loc_ref, kmp_int32 gtid,
+                                  kmp_int32 flags, size_t sizeof_kmp_task_t,
+                                  size_t sizeof_shareds,
+                                  kmp_routine_entry_t task_entry)
+    __attribute__((weak));
+
+kmp_task_t *
+__kmpc_omp_target_task_alloc(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 flags,
+                             size_t sizeof_kmp_task_t, size_t sizeof_shareds,
+                             kmp_routine_entry_t task_entry,
+                             kmp_int64 device_id) __attribute__((weak));
+
+kmp_int32 __kmpc_omp_task_with_deps(ident_t *loc_ref, kmp_int32 gtid,
+                                    kmp_task_t *new_task, kmp_int32 ndeps,
+                                    kmp_depend_info_t *dep_list,
+                                    kmp_int32 ndeps_noalias,
+                                    kmp_depend_info_t *noalias_dep_list)
+    __attribute__((weak));
+
+/**
+ * The argument set that is passed from asynchronous memory copy to block
+ * version of memory copy invoked in helper task
+ */
+struct TargetMemcpyArgsTy {
+  /**
+   * Common attribuutes
+   */
+  void *Dst;
+  const void *Src;
+  int DstDevice;
+  int SrcDevice;
+
+  /**
+   * The flag that denotes single dimensional or rectangle dimensional copy
+   */
+  bool IsRectMemcpy;
+
+  /**
+   * Arguments for single dimensional copy
+   */
+  size_t Length;
+  size_t DstOffset;
+  size_t SrcOffset;
+
+  /**
+   * Arguments for rectangle dimensional copy
+   */
+  size_t ElementSize;
+  int NumDims;
+  const size_t *Volume;
+  const size_t *DstOffsets;
+  const size_t *SrcOffsets;
+  const size_t *DstDimensions;
+  const size_t *SrcDimensions;
+
+  /**
+   * Constructor for single dimensional copy
+   */
+  TargetMemcpyArgsTy(void *Dst, const void *Src, size_t Length,
+                     size_t DstOffset, size_t SrcOffset, int DstDevice,
+                     int SrcDevice)
+      : Dst(Dst), Src(Src), DstDevice(DstDevice), SrcDevice(SrcDevice),
+        IsRectMemcpy(false), Length(Length), DstOffset(DstOffset),
+        SrcOffset(SrcOffset), ElementSize(0), NumDims(0), Volume(0),
+        DstOffsets(0), SrcOffsets(0), DstDimensions(0), SrcDimensions(0){};
+
+  /**
+   * Constructor for rectangle dimensional copy
+   */
+  TargetMemcpyArgsTy(void *Dst, const void *Src, size_t ElementSize,
+                     int NumDims, const size_t *Volume,
+                     const size_t *DstOffsets, const size_t *SrcOffsets,
+                     const size_t *DstDimensions, const size_t *SrcDimensions,
+                     int DstDevice, int SrcDevice)
+      : Dst(Dst), Src(Src), DstDevice(DstDevice), SrcDevice(SrcDevice),
+        IsRectMemcpy(true), Length(0), DstOffset(0), SrcOffset(0),
+        ElementSize(ElementSize), NumDims(NumDims), Volume(Volume),
+        DstOffsets(DstOffsets), SrcOffsets(SrcOffsets),
+        DstDimensions(DstDimensions), SrcDimensions(SrcDimensions){};
+};
 // Invalid GTID as defined by libomp; keep in sync
 #define KMP_GTID_DNE (-2)
 #ifdef __cplusplus
 }
 #endif
 
 #define TARGET_NAME Libomptarget
 #define DEBUG_PREFIX GETNAME(TARGET_NAME)
 
 ////////////////////////////////////////////////////////////////////////////////
 /// dump a table of all the host-target pointer pairs on failure
 static inline void dumpTargetPointerMappings(const ident_t *Loc,
                                              DeviceTy &Device) {
   DeviceTy::HDTTMapAccessorTy HDTTMap =

openmp/libomptarget/test/api/omp_target_memcpy_async1.c

+// RUN: %libomptarget-compile-and-run-generic
+
+// Test case for omp_target_memcpy_async, oringally from GCC
+
+#include "stdio.h"
+#include <omp.h>
+#include <stdlib.h>
+
+int main() {
+  int d = omp_get_default_device();
+  int id = omp_get_initial_device();
+  int q[128], i;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices())
+    d = id;
+
+  p = omp_target_alloc(130 * sizeof(int), d);
+  if (p == NULL)
+    return 0;
+
+  for (i = 0; i < 128; i++)
+    q[i] = i;
+
+  if (omp_target_memcpy_async(p, q, 128 * sizeof(int), sizeof(int), 0, d, id, 0,
+                              NULL)) {
+    abort();
+  }
+
+#pragma omp taskwait
+
+  int q2[128];
+  for (i = 0; i < 128; ++i)
+    q2[i] = 0;
+  if (omp_target_memcpy_async(q2, p, 128 * sizeof(int), 0, sizeof(int), id, d,
+                              0, NULL))
+    abort();
+
+#pragma omp taskwait
+
+  for (i = 0; i < 128; ++i)
+    if (q2[i] != q[i])
+      abort();
+
+  omp_target_free(p, d);
+
+  return 0;
+}

openmp/libomptarget/test/api/omp_target_memcpy_async2.c

+// RUN: %libomptarget-compile-and-run-generic
+
+#include "stdio.h"
+#include <omp.h>
+#include <stdlib.h>
+
+int main() {
+  int d = omp_get_default_device();
+  int id = omp_get_initial_device();
+  int a[128], b[64], c[32], e[16], q[128], i;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices())
+    d = id;
+
+  p = omp_target_alloc(130 * sizeof(int), d);
+  if (p == NULL)
+    return 0;
+
+  for (i = 0; i < 128; ++i)
+    a[i] = i + 1;
+  for (i = 0; i < 64; ++i)
+    b[i] = i + 2;
+  for (i = 0; i < 32; i++)
+    c[i] = 0;
+  for (i = 0; i < 16; i++)
+    e[i] = i + 4;
+
+  omp_depend_t obj[2];
+
+#pragma omp parallel num_threads(5)
+#pragma omp single
+  {
+#pragma omp task depend(out : p)
+    omp_target_memcpy(p, a, 128 * sizeof(int), 0, 0, d, id);
+
+#pragma omp task depend(inout : p)
+    omp_target_memcpy(p, b, 64 * sizeof(int), 0, 0, d, id);
+
+#pragma omp task depend(out : c)
+    for (i = 0; i < 32; i++)
+      c[i] = i + 3;
+
+#pragma omp depobj(obj[0]) depend(inout : p)
+#pragma omp depobj(obj[1]) depend(in : c)
+    omp_target_memcpy_async(p, c, 32 * sizeof(int), 0, 0, d, id, 2, obj);
+
+#pragma omp task depend(in : p)
+    omp_target_memcpy(p, e, 16 * sizeof(int), 0, 0, d, id);
+  }
+
+#pragma omp taskwait
+
+  for (i = 0; i < 128; ++i)
+    q[i] = 0;
+  omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d);
+  for (i = 0; i < 16; ++i)
+    if (q[i] != i + 4)
+      abort();
+  for (i = 16; i < 32; ++i)
+    if (q[i] != i + 3)
+      abort();
+  for (i = 32; i < 64; ++i)
+    if (q[i] != i + 2)
+      abort();
+  for (i = 64; i < 128; ++i)
+    if (q[i] != i + 1)
+      abort();
+
+  omp_target_free(p, d);
+
+  return 0;
+}

openmp/libomptarget/test/api/omp_target_memcpy_rect_async1.c

+// RUN: %libomptarget-compile-and-run-generic
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define NUM_DIMS 3
+
+int main() {
+  int d = omp_get_default_device();
+  int id = omp_get_initial_device();
+  int q[128], q2[128], i;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices())
+    d = id;
+
+  p = omp_target_alloc(130 * sizeof(int), d);
+  if (p == NULL)
+    return 0;
+
+  if (omp_target_memcpy_rect_async(NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
+                                   NULL, d, id, 0, NULL) < 3 ||
+      omp_target_memcpy_rect_async(NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
+                                   NULL, id, d, 0, NULL) < 3 ||
+      omp_target_memcpy_rect_async(NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
+                                   NULL, id, id, 0, NULL) < 3)
+    abort();
+
+  for (i = 0; i < 128; i++)
+    q[i] = 0;
+  if (omp_target_memcpy(p, q, 128 * sizeof(int), 0, 0, d, id) != 0)
+    abort();
+
+  for (i = 0; i < 128; i++)
+    q[i] = i + 1;
+
+  size_t volume[NUM_DIMS] = {1, 2, 3};
+  size_t dst_offsets[NUM_DIMS] = {0, 0, 0};
+  size_t src_offsets[NUM_DIMS] = {0, 0, 0};
+  size_t dst_dimensions[NUM_DIMS] = {3, 4, 5};
+  size_t src_dimensions[NUM_DIMS] = {2, 3, 4};
+
+  if (omp_target_memcpy_rect_async(p, q, sizeof(int), NUM_DIMS, volume,
+                                   dst_offsets, src_offsets, dst_dimensions,
+                                   src_dimensions, d, id, 0, NULL) != 0)
+    abort();
+
+#pragma omp taskwait
+
+  for (i = 0; i < 128; i++)
+    q2[i] = 0;
+  if (omp_target_memcpy(q2, p, 128 * sizeof(int), 0, 0, id, d) != 0)
+    abort();
+
+  /* q2 is expected to contain: 1 2 3 0 0 5 6 7 0 0 .. 0  */
+  if (q2[0] != 1 || q2[1] != 2 || q2[2] != 3 || q2[3] != 0 || q2[4] != 0 ||
+      q2[5] != 5 || q2[6] != 6 || q2[7] != 7)
+    abort();
+  for (i = 8; i < 128; ++i)
+    if (q2[i] != 0)
+      abort();
+
+  omp_target_free(p, d);
+  return 0;
+}

openmp/libomptarget/test/api/omp_target_memcpy_rect_async2.c

+// RUN: %libomptarget-compile-and-run-generic
+
+#include <omp.h>
+#include <stdlib.h>
+
+#define NUM_DIMS 3
+
+int main() {
+  int d = omp_get_default_device();
+  int id = omp_get_initial_device();
+  int a[128], b[64], c[128], e[16], q[128], i;
+  void *p;
+
+  if (d < 0 || d >= omp_get_num_devices())
+    d = id;
+
+  p = omp_target_alloc(130 * sizeof(int), d);
+  if (p == NULL)
+    return 0;
+
+  for (i = 0; i < 128; i++)
+    q[i] = 0;
+  if (omp_target_memcpy(p, q, 128 * sizeof(int), 0, 0, d, id) != 0)
+    abort();
+
+  size_t volume[NUM_DIMS] = {2, 2, 3};
+  size_t dst_offsets[NUM_DIMS] = {0, 0, 0};
+  size_t src_offsets[NUM_DIMS] = {0, 0, 0};
+  size_t dst_dimensions[NUM_DIMS] = {3, 4, 5};
+  size_t src_dimensions[NUM_DIMS] = {2, 3, 4};
+
+  for (i = 0; i < 128; i++)
+    a[i] = 42;
+  for (i = 0; i < 64; i++)
+    b[i] = 24;
+  for (i = 0; i < 128; i++)
+    c[i] = 0;
+  for (i = 0; i < 16; i++)
+    e[i] = 77;
+
+  omp_depend_t obj[2];
+
+#pragma omp parallel num_threads(5)
+#pragma omp single
+  {
+#pragma omp task depend(out : p)
+    omp_target_memcpy(p, a, 128 * sizeof(int), 0, 0, d, id);
+
+#pragma omp task depend(inout : p)
+    omp_target_memcpy(p, b, 64 * sizeof(int), 0, 0, d, id);
+
+#pragma omp task depend(out : c)
+    for (i = 0; i < 128; i++)
+      c[i] = i + 1;
+
+#pragma omp depobj(obj[0]) depend(inout : p)
+#pragma omp depobj(obj[1]) depend(in : c)
+
+    /*  This produces: 1 2 3 - - 5 6 7 - - at positions 0..9 and
+        13 14 15 - - 17 18 19 - - at positions 20..29.  */
+    omp_target_memcpy_rect_async(p, c, sizeof(int), NUM_DIMS, volume,
+                                 dst_offsets, src_offsets, dst_dimensions,
+                                 src_dimensions, d, id, 2, obj);
+
+#pragma omp task depend(in : p)
+    omp_target_memcpy(p, e, 16 * sizeof(int), 0, 0, d, id);
+  }
+
+#pragma omp taskwait
+
+  if (omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d) != 0)
+    abort();
+
+  for (i = 0; i < 16; ++i)
+    if (q[i] != 77)
+      abort();
+  if (q[20] != 13 || q[21] != 14 || q[22] != 15 || q[25] != 17 || q[26] != 18 ||
+      q[27] != 19)
+    abort();
+  for (i = 28; i < 64; ++i)
+    if (q[i] != 24)
+      abort();
+  for (i = 64; i < 128; ++i)
+    if (q[i] != 42)
+      abort();
+
+  omp_target_free(p, d);
+  return 0;
+}