[libomptarget] Add support for target update non-contiguous

[cchen]

Summary

This patch is the runtime support for https://reviews.llvm.org/D84192.

In order not to modify the tgt_target_data_update information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload arg when
the maptype is set as OMP_TGT_MAPTYPE_DESCRIPTOR. The origin arg is for
passing the pointer information, however, the overloaded arg is an
array of descriptor_dim:

struct descriptor_dim {
  int64_t offset;
  int64_t count;
  int64_t stride
};

and the array size is the dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
arg_size parameter by using dimension size.

[ABataev]

The test?

openmp/libomptarget/include/omptarget.h

131–133

clang part relies on unsigned types.

[cchen]

Add test

[ABataev]

openmp/libomptarget/test/offloading/non_contiguous_update.c

1 ↗

(On Diff #272529)

We don't have support for this in clang yet, better to add a test that uses runtime calls directly for the new features.

[cchen]

openmp/libomptarget/test/offloading/non_contiguous_update.c

1 ↗

(On Diff #272529)

Is there any test file that I can refer to? The run line I wrote here is basically copied from other test files under libomptarget/test.

[ABataev]

openmp/libomptarget/test/offloading/non_contiguous_update.c

1 ↗

(On Diff #272529)

Sure, check llvm-project/openmp/libomptarget/test/mapping/declare_mapper_api.cpp for example

[cchen]

Update test

[cchen]

Revert llvm pre-merge lint for existing code

[cchen]

Update test so that it don't depends on compiler generated code

[cchen]

Fix non_contig type

[cchen]

ping

[cchen]

ping

[cchen]

Modify test for supporting stride in target update non-contiguous

[cchen]

Ping, compiler support is in this patch: https://reviews.llvm.org/D84192.

[cchen]

ping

[grokos]

openmp/libomptarget/include/omptarget.h

54

Maybe use a more descriptive name for the flag? Like OMP_TGT_MAPTYPE_NON_CONTIG?

openmp/libomptarget/src/omptarget.cpp

801

Maybe rename dim to current_dim - if it wasn't for the comment in the function call in line 629 I wouldn't have figured out what this argument describes.

809

alreay --> already
peice --> piece

814

return --> returns

[cchen]

Rebase and fix based on feedback

[grokos]

The link in the description for the clang patch is outdated (that patch has been abandoned), can you replace it with https://reviews.llvm.org/D84192? Thanks!

[cchen]

In D82245#2377081, @grokos wrote:

The link in the description for the clang patch is outdated (that patch has been abandoned), can you replace it with https://reviews.llvm.org/D84192? Thanks!

I just updated the link, thanks!

[grokos]

LGTM. Thanks!

[tianshilei1992]

This patch breaks the building of OpenMP. Please fix it.

FAILED: libomptarget/src/CMakeFiles/omptarget.dir/omptarget.cpp.o
/home/shiltian/Documents/deploy/llvm/release/bin/clang++  -DOMPTARGET_DEBUG -Domptarget_EXPORTS -I/home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/include -Wall -Wcast-qual -Wformat-pedantic -Wimplicit-fallthrough -Wsign-compare -Wno-extra -Wno-pedantic -std=gnu++14 -g -fPIC -MD -MT libomptarget/src/CMakeFiles/omptarget.dir/omptarget.cpp.o -MF libomptarget/src/CMakeFiles/omptarget.dir/omptarget.cpp.o.d -o libomptarget/src/CMakeFiles/omptarget.dir/omptarget.cpp.o -c /home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/src/omptarget.cpp
/home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/src/omptarget.cpp:664:8: error: use of undeclared identifier 'arg_sizes'
       arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
       ^
/home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/src/omptarget.cpp:664:18: error: use of undeclared identifier 'i'
       arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
                 ^
/home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/src/omptarget.cpp:691:8: error: use of undeclared identifier 'arg_sizes'
       arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
       ^
/home/shiltian/Documents/vscode/llvm-project/openmp/libomptarget/src/omptarget.cpp:691:18: error: use of undeclared identifier 'i'
       arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
                 ^
4 errors generated.

What's more, please follow LLVM coding style.

openmp/libomptarget/src/omptarget.cpp

668

arg_size is not defined here.

[cchen]

Fix build error

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

642

Since you're modifying this function, could you please change the coding style to LLVM style as well?

[cchen]

openmp/libomptarget/src/omptarget.cpp

642

By coding style do you mean the name of the parameters? Can you be more specific? Thanks.

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

642

Yes, sort of. You can refer to https://llvm.org/docs/CodingStandards.html for more details.

[cchen]

Fix coding style

[protze.joachim]

Please revert your commit while you work on a fix, as it breaks building the runtime code.

[tianshilei1992]

In D82245#2383337, @cchen wrote:

Fix coding style

I'm not sure whether this is still WIP because I didn't see any change in the touched code.

[cchen]

In D82245#2386544, @tianshilei1992 wrote:

In D82245#2383337, @cchen wrote:

Fix coding style

I'm not sure whether this is still WIP because I didn't see any change in the touched code.

I've changed the name of the argument in that function and apply clang-format on the function. Is there any place that I'm not following the coding style? I'm just waiting for acceptance to merge the patch.

[tianshilei1992]

In D82245#2386727, @cchen wrote:

In D82245#2386544, @tianshilei1992 wrote:

In D82245#2383337, @cchen wrote:

Fix coding style

I'm not sure whether this is still WIP because I didn't see any change in the touched code.

I've changed the name of the argument in that function and apply clang-format on the function. Is there any place that I'm not following the coding style? I'm just waiting for acceptance to merge the patch.

Take the inline comment as an example, but for those C functions, just leave them what they're now.

openmp/libomptarget/src/omptarget.cpp

642

Based on LLVM standard, this function should be

static int targetDataContiguous(DeviceTy &Device, void *ArgsBase, void *Arg, int64_t argSize, int64_t ArgType)

642

static int targetDataContiguous(DeviceTy &Device, void *ArgsBase, void *Arg, int64_t ArgSize, int64_t ArgType)

[cchen]

Apply LLVM coding style

[tianshilei1992]

openmp/libomptarget/include/omptarget.h

130

Like the void *Queue in __tgt_async_info

openmp/libomptarget/src/omptarget.cpp

669

Still not LLVM style.

675

ditto

676

ditto

678

ditto

[cchen]

Update style

[tianshilei1992]

Would you please take a thorough check with your patch? I can still find many code that still don't conform with the standard. Thanks. The test case is fine as we don't have unified standard.

openmp/libomptarget/src/omptarget.cpp

641–642

This function does not conform with the standard.

[cchen]

In D82245#2404499, @tianshilei1992 wrote:

Would you please take a thorough check with your patch? I can still find many code that still don't conform with the standard. Thanks. The test case is fine as we don't have unified standard.

Sure, I'm just not sure when to follow the style since the style in this codebase is so inconsistent and it confused if I should modify code not written by me.

[cchen]

Refactor target_data_update

[tianshilei1992]

LGTM. Thanks.

openmp/libomptarget/include/omptarget.h

enum tgt_map_type {
   OMP_TGT_MAPTYPE_PRIVATE         = 0x080,
   // copy by value - not mapped
   OMP_TGT_MAPTYPE_LITERAL         = 0x100,
   // mapping is implicit
   OMP_TGT_MAPTYPE_IMPLICIT        = 0x200,
   // copy data to device
   OMP_TGT_MAPTYPE_CLOSE           = 0x400,
   // runtime error if not already allocated
   OMP_TGT_MAPTYPE_PRESENT         = 0x1000,
+  // descriptor for non-contiguous target-update
+  OMP_TGT_MAPTYPE_NON_CONTIG      = 0x100000000000,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };
 
 enum OpenMPOffloadingDeclareTargetFlags {
   /// Mark the entry as having a 'link' attribute.
   OMP_DECLARE_TARGET_LINK = 0x01,
   /// Mark the entry as being a global constructor.

 /// operations for device-dependent optimization and potential synchronization
 struct __tgt_async_info {
   // A pointer to a queue-like structure where offloading operations are issued.
   // We assume to use this structure to do synchronization. In CUDA backend, it
   // is CUstream.
   void *Queue = nullptr;
 };
 
+/// This struct is a record of non-contiguous information
+struct __tgt_target_non_contig {
+  uint64_t Offset;
+  uint64_t Count;
+  uint64_t Stride;
+};
+
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 int omp_get_num_devices(void);
 int omp_get_initial_device(void);
 void *omp_target_alloc(size_t size, int device_num);
 void omp_target_free(void *device_ptr, int device_num);

openmp/libomptarget/src/device.cpp

HostDataToTargetMap.emplace(
                            (uintptr_t)HstPtrBegin + Size, tp, HstPtrName));
     rc = (void *)tp;
   }
 
   DataMapMtx.unlock();
   return rc;
 }
 
-// Used by targetDataBegin, targetDataEnd, target_data_update and target.
+// Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
                                bool UpdateRefCount, bool &IsHostPtr,
                                bool MustContain) {
   void *rc = NULL;
   IsHostPtr = false;
   IsLast = false;

openmp/libomptarget/src/interface.cpp

EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     DP("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false, loc);
     return;
   }
 
   DeviceTy &Device = PM->Devices[device_id];
-  int rc = target_data_update(Device, arg_num, args_base, args, arg_sizes,
-                              arg_types, arg_names, arg_mappers);
+  int rc = targetDataUpdate(Device, arg_num, args_base, args, arg_sizes,
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
 }
 
 EXTERN void __tgt_target_data_update_nowait_mapper(
     ident_t *loc, int64_t device_id, int32_t arg_num, void **args_base,
     void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {

openmp/libomptarget/src/omptarget.cpp

for (DeallocTgtPtrInfo &Info : DeallocTgtPtrs) {
     if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Deallocating data from device failed.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
-/// Internal function to pass data to/from the target.
-// async_info_ptr is currently unused, added here so target_data_update has the
-// same signature as targetDataBegin and targetDataEnd.
-int target_data_update(DeviceTy &Device, int32_t arg_num, void **args_base,
-                       void **args, int64_t *arg_sizes, int64_t *arg_types,
-                       map_var_info_t *arg_names, void **arg_mappers,
-                       __tgt_async_info *async_info_ptr) {
-  // process each input.
-  for (int32_t i = 0; i < arg_num; ++i) {
-    if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) ||
-        (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE))
-      continue;
-
-    if (arg_mappers && arg_mappers[i]) {
-      // Instead of executing the regular path of target_data_update, call the
-      // targetDataMapper variant which will call target_data_update again
-      // with new arguments.
-      DP("Calling targetDataMapper for the %dth argument\n", i);
-
-      int rc =
-          targetDataMapper(Device, args_base[i], args[i], arg_sizes[i],
-                           arg_types[i], arg_mappers[i], target_data_update);
-
-      if (rc != OFFLOAD_SUCCESS) {
-        REPORT(
-            "Call to target_data_update via targetDataMapper for custom mapper"
-            " failed.\n");
-        return OFFLOAD_FAIL;
-      }
-
-      // Skip the rest of this function, continue to the next argument.
-      continue;
-    }
-
-    void *HstPtrBegin = args[i];
-    int64_t MapSize = arg_sizes[i];
-    bool IsLast, IsHostPtr;
-    void *TgtPtrBegin = Device.getTgtPtrBegin(
-        HstPtrBegin, MapSize, IsLast, false, IsHostPtr, /*MustContain=*/true);
-    if (!TgtPtrBegin) {
-      DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
-      if (arg_types[i] & OMP_TGT_MAPTYPE_PRESENT) {
-        MESSAGE("device mapping required by 'present' motion modifier does not "
-                "exist for host address " DPxMOD " (%" PRId64 " bytes)",
-                DPxPTR(HstPtrBegin), MapSize);
-        return OFFLOAD_FAIL;
-      }
-      continue;
-    }
+static int targetDataContiguous(DeviceTy &Device, void *ArgsBase,
+                                void *HstPtrBegin, int64_t ArgSize,
+                                int64_t ArgType) {
+  void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSize, IsLast, false,
+                                            IsHostPtr, /*MustContain=*/true);
+    if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
+              DPxPTR(HstPtrBegin), ArgSize);
+    return OFFLOAD_SUCCESS;
 
-    if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
-        TgtPtrBegin == HstPtrBegin) {
-      DP("hst data:" DPxMOD " unified and shared, becomes a noop\n",
-         DPxPTR(HstPtrBegin));
-      continue;
-    }
+    return OFFLOAD_SUCCESS;
 
-    if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
-      DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
-          arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
-      int rt = Device.retrieveData(HstPtrBegin, TgtPtrBegin, MapSize, nullptr);
-      if (rt != OFFLOAD_SUCCESS) {
-        REPORT("Copying data from device failed.\n");
-        return OFFLOAD_FAIL;
-      }
+  if (ArgType & OMP_TGT_MAPTYPE_FROM) {
+       ArgSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+    int Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, ArgSize, nullptr);
+    if (Ret != OFFLOAD_SUCCESS) {
 
-      uintptr_t lb = (uintptr_t) HstPtrBegin;
-      uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
-      Device.ShadowMtx.lock();
-      for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
-          it != Device.ShadowPtrMap.end(); ++it) {
-        void **ShadowHstPtrAddr = (void**) it->first;
-        if ((uintptr_t) ShadowHstPtrAddr < lb)
-          continue;
-        if ((uintptr_t) ShadowHstPtrAddr >= ub)
-          break;
-        DP("Restoring original host pointer value " DPxMOD " for host pointer "
-            DPxMOD "\n", DPxPTR(it->second.HstPtrVal),
-            DPxPTR(ShadowHstPtrAddr));
-        *ShadowHstPtrAddr = it->second.HstPtrVal;
-      }
-      Device.ShadowMtx.unlock();
-    }
+    uintptr_t LB = (uintptr_t)HstPtrBegin;
+    uintptr_t UB = (uintptr_t)HstPtrBegin + ArgSize;
+    for (ShadowPtrListTy::iterator IT = Device.ShadowPtrMap.begin();
+         IT != Device.ShadowPtrMap.end(); ++IT) {
+      void **ShadowHstPtrAddr = (void **)IT->first;
+      if ((uintptr_t)ShadowHstPtrAddr < LB)
+      if ((uintptr_t)ShadowHstPtrAddr >= UB)
+      DP("Restoring original host pointer value " DPxMOD
+         " for host pointer " DPxMOD "\n",
+         DPxPTR(IT->second.HstPtrVal), DPxPTR(ShadowHstPtrAddr));
+      *ShadowHstPtrAddr = IT->second.HstPtrVal;
 
-    if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
-      DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
-          arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
-      int rt = Device.submitData(TgtPtrBegin, HstPtrBegin, MapSize, nullptr);
-      if (rt != OFFLOAD_SUCCESS) {
-        REPORT("Copying data to device failed.\n");
-        return OFFLOAD_FAIL;
-      }
+  if (ArgType & OMP_TGT_MAPTYPE_TO) {
+       ArgSize, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
+    int Ret = Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSize, nullptr);
+    if (Ret != OFFLOAD_SUCCESS) {
 
-      uintptr_t lb = (uintptr_t) HstPtrBegin;
-      uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
-      Device.ShadowMtx.lock();
-      for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
-          it != Device.ShadowPtrMap.end(); ++it) {
-        void **ShadowHstPtrAddr = (void **)it->first;
-        if ((uintptr_t)ShadowHstPtrAddr < lb)
-          continue;
-        if ((uintptr_t)ShadowHstPtrAddr >= ub)
-          break;
-        DP("Restoring original target pointer value " DPxMOD " for target "
-           "pointer " DPxMOD "\n",
-           DPxPTR(it->second.TgtPtrVal), DPxPTR(it->second.TgtPtrAddr));
-        rt = Device.submitData(it->second.TgtPtrAddr, &it->second.TgtPtrVal,
-                               sizeof(void *), nullptr);
-        if (rt != OFFLOAD_SUCCESS) {
-          REPORT("Copying data to device failed.\n");
-          Device.ShadowMtx.unlock();
-          return OFFLOAD_FAIL;
-        }
-      }
-      Device.ShadowMtx.unlock();
-    }
+    uintptr_t LB = (uintptr_t)HstPtrBegin;
+    uintptr_t UB = (uintptr_t)HstPtrBegin + ArgSize;
+    for (ShadowPtrListTy::iterator IT = Device.ShadowPtrMap.begin();
+         IT != Device.ShadowPtrMap.end(); ++IT) {
+      void **ShadowHstPtrAddr = (void **)IT->first;
+      if ((uintptr_t)ShadowHstPtrAddr < LB)
+      if ((uintptr_t)ShadowHstPtrAddr >= UB)
+         DPxPTR(IT->second.TgtPtrVal), DPxPTR(IT->second.TgtPtrAddr));
+      Ret = Device.submitData(IT->second.TgtPtrAddr, &IT->second.TgtPtrVal,
+      if (Ret != OFFLOAD_SUCCESS) {
+  return OFFLOAD_SUCCESS;
+}
+
+static int targetDataNonContiguous(DeviceTy &Device, void *ArgsBase,
+                                   __tgt_target_non_contig *NonContig,
+                                   uint64_t Size, int64_t ArgType,
+                                   int CurrentDim, int DimSize,
+                                   uint64_t Offset) {
+  int Ret = OFFLOAD_SUCCESS;
+  if (CurrentDim < DimSize) {
+    for (unsigned int I = 0; I < NonContig[CurrentDim].Count; ++I) {
+      uint64_t CurOffset =
+          (NonContig[CurrentDim].Offset + I) * NonContig[CurrentDim].Stride;
+      // we only need to transfer the first element for the last dimension
+      // since we've already got a contiguous piece.
+      if (CurrentDim != DimSize - 1 || I == 0) {
+        Ret = targetDataNonContiguous(Device, ArgsBase, NonContig, Size,
+                                      ArgType, CurrentDim + 1, DimSize,
+                                      Offset + CurOffset);
+        // Stop the whole process if any contiguous piece returns anything
+        // other than OFFLOAD_SUCCESS.
+        if (Ret != OFFLOAD_SUCCESS)
+          return Ret;
+      }
+    }
+  } else {
+    char *Ptr = (char *)ArgsBase + Offset;
+    DP("Transfer of non-contiguous : host ptr %lx offset %ld len %ld\n",
+       (uint64_t)Ptr, Offset, Size);
+    Ret = targetDataContiguous(Device, ArgsBase, Ptr, Size, ArgType);
+  }
+  return Ret;
+}
+
+static int getNonContigMergedDimension(__tgt_target_non_contig *NonContig,
+                                       int32_t DimSize) {
+  int RemovedDim = 0;
+  for (int I = DimSize - 1; I > 0; --I) {
+    if (NonContig[I].Count * NonContig[I].Stride == NonContig[I - 1].Stride)
+      RemovedDim++;
+  }
+  return RemovedDim;
+}
+
+/// Internal function to pass data to/from the target.
+// async_info_ptr is currently unused, added here so targetDataUpdate has the
+// same signature as targetDataBegin and targetDataEnd.
+int targetDataUpdate(DeviceTy &Device, int32_t ArgNum, void **ArgsBase,
+                     void **Args, int64_t *ArgSizes, int64_t *ArgTypes,
+                     map_var_info_t *ArgNames, void **ArgMappers,
+                     __tgt_async_info *AsyncInfoPtr) {
+  // process each input.
+  for (int32_t I = 0; I < ArgNum; ++I) {
+    if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
+        (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
+      continue;
+
+    if (ArgMappers && ArgMappers[I]) {
+      // Instead of executing the regular path of targetDataUpdate, call the
+      // targetDataMapper variant which will call targetDataUpdate again
+      // with new arguments.
+      DP("Calling targetDataMapper for the %dth argument\n", I);
+
+      int Ret = targetDataMapper(Device, ArgsBase[I], Args[I], ArgSizes[I],
+                                 ArgTypes[I], ArgMappers[I], targetDataUpdate);
+
+      if (Ret != OFFLOAD_SUCCESS) {
+        REPORT("Call to targetDataUpdate via targetDataMapper for custom mapper"
+               " failed.\n");
+        return OFFLOAD_FAIL;
+      }
+
+      // Skip the rest of this function, continue to the next argument.
+      continue;
+    }
+
+    int Ret = OFFLOAD_SUCCESS;
+
+    if (ArgTypes[I] & OMP_TGT_MAPTYPE_NON_CONTIG) {
+      __tgt_target_non_contig *NonContig = (__tgt_target_non_contig *)Args[I];
+      int32_t DimSize = ArgSizes[I];
+      uint64_t Size =
+          NonContig[DimSize - 1].Count * NonContig[DimSize - 1].Stride;
+      int32_t MergedDim = getNonContigMergedDimension(NonContig, DimSize);
+      Ret = targetDataNonContiguous(
+          Device, ArgsBase[I], NonContig, Size, ArgTypes[I],
+          /*current_dim=*/0, DimSize - MergedDim, /*offset=*/0);
+    } else {
+      Ret = targetDataContiguous(Device, ArgsBase[I], Args[I], ArgSizes[I],
+                                 ArgTypes[I]);
+    }
+    if (Ret == OFFLOAD_FAIL)
+      return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }
 
 static const unsigned LambdaMapping = OMP_TGT_MAPTYPE_PTR_AND_OBJ |
                                       OMP_TGT_MAPTYPE_LITERAL |
                                       OMP_TGT_MAPTYPE_IMPLICIT;
 static bool isLambdaMapping(int64_t Mapping) {

openmp/libomptarget/src/private.h

extern int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
                            map_var_info_t *arg_names, void **arg_mappers,
                            __tgt_async_info *async_info_ptr);
 
 extern int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
                          void **Args, int64_t *ArgSizes, int64_t *ArgTypes,
                          map_var_info_t *arg_names, void **ArgMappers,
                          __tgt_async_info *AsyncInfo);
 
-extern int target_data_update(DeviceTy &Device, int32_t arg_num,
-                              void **args_base, void **args, int64_t *arg_sizes,
-                              int64_t *arg_types, map_var_info_t *arg_names,
-                              void **arg_mappers,
-                              __tgt_async_info *async_info_ptr = nullptr);
+extern int targetDataUpdate(DeviceTy &Device, int32_t arg_num, void **args_base,
+                            void **args, int64_t *arg_sizes, int64_t *arg_types,
+                            map_var_info_t *arg_names, void **arg_mappers,
 
 extern int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, map_var_info_t *arg_names,
                   void **ArgMappers, int32_t TeamNum, int32_t ThreadLimit,
                   int IsTeamConstruct);
 
 extern int CheckDeviceAndCtors(int64_t device_id);

 
 // The mapper function pointer type. It follows the signature below:
 // void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 //                                           void *base, void *begin,
 //                                           size_t size, int64_t type);
 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
 
 // Function pointer type for target_data_* functions (targetDataBegin,
-// targetDataEnd and target_data_update).
+// targetDataEnd and targetDataUpdate).
 typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
                                    int64_t *, int64_t *, map_var_info_t *,
                                    void **, __tgt_async_info *);
 
 // Implemented in libomp, they are called from within __tgt_* functions.
 #ifdef __cplusplus
 extern "C" {
 #endif

openmp/libomptarget/test/offloading/non_contiguous_update.cpp

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
+
+#include <cstdio>
+#include <cstdlib>
+#include <cassert>
+
+// Data structure definitions copied from OpenMP RTL.
+struct __tgt_target_non_contig {
+  int64_t offset;
+  int64_t width;
+  int64_t stride;
+};
+
+enum tgt_map_type {
+  OMP_TGT_MAPTYPE_NON_CONTIG      = 0x100000000000
+};
+
+// OpenMP RTL interfaces
+#ifdef __cplusplus
+extern "C" {
+#endif
+void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
+                              void **args_base, void **args, int64_t *arg_sizes,
+                              int64_t *arg_types);
+#ifdef __cplusplus
+}
+#endif
+
+int main() {
+  // case 1
+  // int arr[3][4][5][6];
+  // #pragma omp target update to(arr[0:2][1:3][1:2][:])
+  // set up descriptor
+  __tgt_target_non_contig non_contig[5] = {
+      {0, 2, 480}, {1, 3, 120}, {1, 2, 24}, {0, 6, 4}, {0, 1, 4}};
+  int64_t size = 4, type = OMP_TGT_MAPTYPE_NON_CONTIG;
+
+  void *base;
+  void *begin = &non_contig;
+  int64_t *sizes = &size;
+  int64_t *types = &type;
+
+  // The below diagram is the visualization of the non-contiguous transfer after
+  // optimization. Note that each element represent the innermost dimension
+  // (unit size = 24) since the stride * count of last dimension is equal to the
+  // stride of second last dimension.
+  //
+  // OOOOO OOOOO OOOOO
+  // OXXOO OXXOO OOOOO
+  // OXXOO OXXOO OOOOO
+  // OXXOO OXXOO OOOOO
+  __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+                           sizes, types);
+  // DEBUG: offset 144
+  // DEBUG: offset 264
+  // DEBUG: offset 384
+  // DEBUG: offset 624
+  // DEBUG: offset 744
+  // DEBUG: offset 864
+
+
+  // case 2
+  // double darr[3][4][5];
+  // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
+  // set up descriptor
+  __tgt_target_non_contig non_contig_2[4] = {
+      {0, 2, 320}, {2, 2, 40}, {0, 2, 16}, {0, 1, 8}};
+  int64_t size_2 = 4, type_2 = OMP_TGT_MAPTYPE_NON_CONTIG;
+
+  void *base_2;
+  void *begin_2 = &non_contig_2;
+  int64_t *sizes_2 = &size_2;
+  int64_t *types_2 = &type_2;
+
+  // The below diagram is the visualization of the non-contiguous transfer after
+  // optimization. Note that each element represent the innermost dimension
+  // (unit size = 24) since the stride * count of last dimension is equal to the
+  // stride of second last dimension.
+  //
+  // OOOOO OOOOO OOOOO
+  // OOOOO OOOOO OOOOO
+  // XOXOO OOOOO XOXOO
+  // XOXOO OOOOO XOXOO
+  __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2,
+                           sizes_2, types_2);
+  // DEBUG: offset 80
+  // DEBUG: offset 96
+  // DEBUG: offset 120
+  // DEBUG: offset 136
+  // DEBUG: offset 400
+  // DEBUG: offset 416
+  // DEBUG: offset 440
+  // DEBUG: offset 456
+  return 0;
+}
+