[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

129–131

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

52

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

openmp/libomptarget/src/omptarget.cpp

802

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.

810

alreay --> already
peice --> piece

815

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

664

arg_size is not defined here.

[cchen]

Fix build error

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

638

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

[cchen]

openmp/libomptarget/src/omptarget.cpp

638

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

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

638

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

638

Based on LLVM standard, this function should be

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

638

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

128

Like the void *Queue in __tgt_async_info

openmp/libomptarget/src/omptarget.cpp

665

Still not LLVM style.

671

ditto

672

ditto

674

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

768

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

 //===-------- 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

 #define OFFLOAD_SUCCESS (0)
 #define OFFLOAD_FAIL (~0)
 
 #define OFFLOAD_DEVICE_DEFAULT     -1
 
 /// Data attributes for each data reference used in an OpenMP target region.
 enum tgt_map_type {
   // No flags
   OMP_TGT_MAPTYPE_NONE            = 0x000,
   // copy data from host to device
   OMP_TGT_MAPTYPE_TO              = 0x001,
   // copy data from device to host
   OMP_TGT_MAPTYPE_FROM            = 0x002,
   // copy regardless of the reference count
   OMP_TGT_MAPTYPE_ALWAYS          = 0x004,
   // force unmapping of data
   OMP_TGT_MAPTYPE_DELETE          = 0x008,
   // map the pointer as well as the pointee
   OMP_TGT_MAPTYPE_PTR_AND_OBJ     = 0x010,
   // pass device base address to kernel
   OMP_TGT_MAPTYPE_TARGET_PARAM    = 0x020,
   // return base device address of mapped data
   OMP_TGT_MAPTYPE_RETURN_PARAM    = 0x040,
   // private variable - not mapped
   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.
   OMP_DECLARE_TARGET_CTOR = 0x02,
   /// Mark the entry as being a global destructor.

 /// 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/omptarget.cpp

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,
-    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 target_data_contiguous(DeviceTy &Device, void *args_base, void *args,
+                                  int64_t arg_size, int64_t arg_type) {
+  void *HstPtrBegin = args;
+      HstPtrBegin, arg_size, IsLast, false, IsHostPtr, /*MustContain=*/true);
+    if (arg_type & OMP_TGT_MAPTYPE_PRESENT) {
+              DPxPTR(HstPtrBegin), arg_size);
+    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 (arg_type & OMP_TGT_MAPTYPE_FROM) {
+       arg_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+    int rt = Device.retrieveData(HstPtrBegin, TgtPtrBegin, arg_size, nullptr);
 
-      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 + arg_size;
+      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));
 
-    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 (arg_type & OMP_TGT_MAPTYPE_TO) {
+       arg_size, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
+    int rt = Device.submitData(TgtPtrBegin, HstPtrBegin, arg_size, nullptr);
 
-      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 + arg_size;
+  return OFFLOAD_SUCCESS;
+}
+
+static int target_data_non_contiguous(DeviceTy &Device, void *arg_base,
+                                      __tgt_target_non_contig *non_contig,
+                                      uint64_t size, int64_t arg_type,
+                                      int current_dim, int dim_size,
+                                      uint64_t offset) {
+  int rt = OFFLOAD_SUCCESS;
+  if (current_dim < dim_size) {
+    for (unsigned int i = 0; i < non_contig[current_dim].count; ++i) {
+      uint64_t cur_offset =
+          (non_contig[current_dim].offset + i) * non_contig[current_dim].stride;
+      // we only need to transfer the first element for the last dimension
+      // since we've already got a contiguous piece.
+      if (current_dim != dim_size - 1 || i == 0) {
+        rt = target_data_non_contiguous(Device, arg_base, non_contig, size,
+                                        arg_type, current_dim + 1, dim_size,
+                                        offset + cur_offset);
+        // Stop the whole process if any contiguous piece returns anything
+        // other than OFFLOAD_SUCCESS.
+        if (rt != OFFLOAD_SUCCESS)
+          return rt;
+      }
+    }
+  } else {
+    char *ptr = (char *)arg_base + offset;
+    DP("Transfer of non-contiguous : host ptr %lx offset %ld len %ld\n",
+       (uint64_t)ptr, offset, size);
+    rt = target_data_contiguous(Device, arg_base, ptr, size, arg_type);
+  }
+  return rt;
+}
+
+static int get_non_contig_merged_dimension(__tgt_target_non_contig *non_contig,
+                                           int32_t dim_size) {
+  int removed_dim = 0;
+  for (int i = dim_size - 1; i > 0; --i) {
+    if (non_contig[i].count * non_contig[i].stride == non_contig[i - 1].stride)
+      removed_dim++;
+  }
+  return removed_dim;
+}
+
+/// 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,
+    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;
+    }
+
+    int rt = OFFLOAD_SUCCESS;
+
+    if (arg_types[i] & OMP_TGT_MAPTYPE_NON_CONTIG) {
+      __tgt_target_non_contig *non_contig = (__tgt_target_non_contig *)args[i];
+      int32_t dim_size = arg_sizes[i];
+      uint64_t size =
+          non_contig[dim_size - 1].count * non_contig[dim_size - 1].stride;
+      int32_t merged_dim =
+          get_non_contig_merged_dimension(non_contig, dim_size);
+      rt = target_data_non_contiguous(
+          Device, args_base[i], non_contig, size, arg_types[i],
+          /*current_dim=*/0, dim_size - merged_dim, /*offset=*/0);
+    } else {
+      rt = target_data_contiguous(Device, args_base[i], args[i], arg_sizes[i],
+                                  arg_types[i]);
+    }
+    if (rt == 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/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;
+}