[OpenMP] Pack first-private arguments to improve efficiency of data transfer

[tianshilei1992]

Summary

In this patch, we pack all small first-private arguments, allocate and transfer them all at once to reduce the number of data transfer which is very expensive.

Let's take the test case as example.

int main() {
  int data1[3] = {1}, data2[3] = {2}, data3[3] = {3};
  int sum[16] = {0};
#pragma omp target teams distribute parallel for map(tofrom: sum) firstprivate(data1, data2, data3)
  for (int i = 0; i < 16; ++i) {
    for (int j = 0; j < 3; ++j) {
      sum[i] += data1[j];
      sum[i] += data2[j];
      sum[i] += data3[j];
    }
  }
}

Here data1, data2, and data3 are three first-private arguments of the target region. In the previous libomptarget, it called data allocation and data transfer three times, each of which allocated and transferred 12 bytes. With this patch, it only calls allocation and transfer once. The size is (12+4)*3=48 where 12 is the size of each array and 4 is the padding to keep the address aligned with 8. It is implemented in this way:

1. First collect all information for those *first*-private arguments. _private_ arguments are not the case because private arguments don't need to be mapped to target device. It just needs a data allocation. With the patch for memory manager, the data allocation could be very cheap, especially for the small size. For each qualified argument, push a place holder pointer nullptr to the vector for kernel arguments, and we will update them later.
2. After we have all information, create a buffer that can accommodate all arguments plus their paddings. Copy the arguments to the buffer at the right place, i.e. aligned address.
3. Allocate a target memory with the same size as the host buffer, transfer the host buffer to target device, and finally update all place holder pointers in the arguments vector.

The reason we only consider small arguments is, the data transfer is asynchronous. Therefore, for the large argument, we could continue to do things on the host side meanwhile, hopefully, the data is also being transferred. The "small" is defined by that the argument size is less than a predefined value. Currently it is 1024. I'm not sure whether it is a good one, and that is an open question. Another question is, do we need to make it configurable via an environment variable?

[tianshilei1992]

This patch contains some changes by clang-format. I'll revert those unrelated changes.

[tianshilei1992]

Removed unrelated changes caused by clang-format

[jdoerfert]

Can you add some C pseudo code to the commit message to illustrate what is happening?

[ye-luo]

Why just "small" ones? why not all of them?

[jdoerfert]

In D86307#2229032, @ye-luo wrote:

Why just "small" ones? why not all of them?

In addition to the last paragraph of the new commit message, we also have to copy the data on the host in the right place. That is not free as the size grows.

[ye-luo]

In D86307#2229059, @jdoerfert wrote:

In D86307#2229032, @ye-luo wrote:

Why just "small" ones? why not all of them?

In addition to the last paragraph of the new commit message, we also have to copy the data on the host in the right place. That is not free as the size grows.

Copy cost on the host is a convincing argument but the asynchronous benefit doesn't seem to me true.

[jdoerfert]

Only minor nits.

openmp/libomptarget/src/omptarget.cpp

841

When I read FP I think floating point. Maybe rename all FPs into FirstPriv or similar, characters are free, confusion is costly.

1017

Why? Still multiple.

1073

when you talk about "the threshold" in the comments above it is not clear what you mean. Add a (see ...) or additional words, e.g., firstprivate bundling threshold.

1104

Make HstPtr a char* in the struct. Then you don't need to unpack hstptr nor size here.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

1065

When I read here with a condition on the transfer.
I'm wondering is this just first-private? or private is also affected?

[JonChesterfield]

This is a good direction. Packing copies together is likely to be faster.

Things like the size limit for profitability will be target specific, suggest adding a hook to the plugin to query such things. Better than yet another environment variable imo.

E.g. over pcie vs a shared memory system are likely to have different thresholds. Nvptx and amdgcn will be different too.

[tianshilei1992]

Update based on comments

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1065

Just first-private because we don't need data transfer for private arguments.

[tianshilei1992]

Updated some comments

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

1065

I'm wondering is this just first-private? or private is also affected?

I were referring to the whole patch instead of just the transfer here. I guess you patch affect both private and firstprivate.
if I have private(data1, data2, data3), will your change make a single allocation?

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1065

No. In that case, MapTo will be false.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

1065

1. Then rename it IsFirstPrivate. MapTo seems quite misleading.
2. In the future, I think it can be nice optimization to allocate all the private in a single shot. Device malloc is still very expensive.

[tianshilei1992]

Rename MapTo to IsFirstPrivate

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1065

1. Done.
2. I think it depends. Just like I said in the code comments somewhere, we have memory manager now, then the allocation might not go to device directly, especially those small size of memory. If we take all private arguments into account, it will increase more host overhead (more host memory, and extra Boolean variable to tell whether we need to copy from original place to the buffer), and also transfer unnecessary data to target device. In fact, in my first implementation, I did have all private arguments, but later I changed it.

[ye-luo]

Only minor things.

openmp/libomptarget/src/omptarget.cpp

830

const?

836

const?

1017

Why? Still multiple.

I think it is a vector of first private arrays. So the name better to be plural.

FirstPrivateArgs seems to be a better name.

[tianshilei1992]

Updated based on comments

[tianshilei1992]

Renamed the vector for first private arguments

[tianshilei1992]

Use synchronous data transfer for the buffer, but I'm not sure whether this can offset the benefit of packing memory copy, even degrade the performance.

[JonChesterfield]

openmp/libomptarget/src/omptarget.cpp

1113

This seems target specific. Multiple async copies can easily be cheaper than one larger sync copy.

How is the lifetime of the host buffer usually managed for async copies? It seems that keeping it alive until the async copy is done would be reasonable, or we could move ownership of the host buffer into the target plugin which is then responsible for the cleanup.

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

1113

I’ll change this part back to async copy by taking the buffer out of the function such that we can make sure that it will be deallocated safely.

[tianshilei1992]

Swtiched to PrivateArgumentManagerTy to handle private arguments

[ye-luo]

Down the road, we may need a way to allocate host pinned memory via the plugin for the host buffer to maximize transfer performance.

openmp/libomptarget/src/omptarget.cpp

872

I think only an index instead of TgtArgs is needed.

924

Document the use of TgtArgs please.

970

In my first thought, I feel better to mark this function private and called by the destructor only. Requiring free() to be called explicitly is error-prone. In a second thought, this probably needed to propagate the return error.

1066

Directly use ArgSizes[I] in the next line.

[tianshilei1992]

Updated based on comments

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

872

It would be not very straightforward why we need an index here, and there is no difference in terms of performance between an index and a reference.

970

Yes, that's exactly the reason that we have a function here.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

872

Without documentation, nothing is straightforward. IndexInTgtArgs is better than TgtArgs. Passing TgtArgs directly open doors to unintended uses.

[tianshilei1992]

Now pass the index of TgtArgs instead of a reference to TgtArgs

[ye-luo]

I prefer to PrivateArgumentManagerTy moved into its own files.
The rest looks good to me.

[ABataev]

openmp/libomptarget/test/mapping/private_mapping.c

2

Why the test is marked as expected to fail? Is this intentional?

[tianshilei1992]

openmp/libomptarget/test/mapping/private_mapping.c

2

That is a mistake. Nice catch, thanks! I'll fix it in another patch. I copied this part from another file and didn't notice it is expected to fail. I did the test manually w/o lit so didn't observe this mistake. I'm wondering why the Buildkite didn't report any issue?

[grokos]

Minor comments (about typos) to be taken into account in case of a future patch.

openmp/libomptarget/src/omptarget.cpp

831

Here you are referring to TgtArgs which is an argument of packAndTransfer(). Can you add a comment in parenthesis to make it clear? The first time I read the code I thought you were mistakenly using the name TgtArgs instead of TgtPtrs (which is a member of PrivateArgumentManagerTy class).

907

needs to mapped --> needs to be copied to target device
Technically speaking, first-private variables are not mapped, they are only copied to the device.

928

is empty --> is not empty

957

by mistaken --> by mistake

openmp/libomptarget/src/omptarget.cpp

if (I != Device.LoopTripCnt.end()) {
       Device.LoopTripCnt.erase(I);
       DP("loop trip count is %lu.\n", LoopTripCount);
     }
   }
 
   return LoopTripCount;
 }
 
+/// A class manages private arguments in a target region.
+class PrivateArgumentManagerTy {
+  /// A data structure for the information of first-private arguments. We can
+  /// use this information to optimize data transfer by packing all
+  /// first-private arguments and transfer them all at once.
+  struct FirstPrivateArgInfoTy {
+    /// The index of the element in \p TgtArgs corresponding to the argument
+    const int Index;
+    /// Host pointer begin
+    const char *HstPtrBegin;
+    /// Host pointer end
+    const char *HstPtrEnd;
+    /// Aligned size
+    const int64_t AlignedSize;
+
+    FirstPrivateArgInfoTy(int Index, const void *HstPtr, int64_t Size)
+        : Index(Index), HstPtrBegin(reinterpret_cast<const char *>(HstPtr)),
+          HstPtrEnd(HstPtrBegin + Size), AlignedSize(Size + Size % Alignment) {}
+  };
+
+  /// A vector of target pointers for all private arguments
+  std::vector<void *> TgtPtrs;
+
+  /// A vector of information of all first-private arguments to be packed
+  std::vector<FirstPrivateArgInfoTy> FirstPrivateArgInfo;
+  /// Host buffer for all arguments to be packed
+  std::vector<char> FirstPrivateArgBuffer;
+  /// The total size of all arguments to be packed
+  int64_t FirstPrivateArgSize = 0;
+
+  /// A reference to the \p DeviceTy object
+  DeviceTy &Device;
+  /// A pointer to a \p __tgt_async_info object
+  __tgt_async_info *AsyncInfo;
+
+  // TODO: What would be the best value here? Should we make it configurable?
+  // If the size is larger than this threshold, we will allocate and transfer it
+  // immediately instead of packing it.
+  static constexpr const int64_t FirstPrivateArgSizeThreshold = 1024;
+
+public:
+  /// Constructor
+  PrivateArgumentManagerTy(DeviceTy &Dev, __tgt_async_info *AsyncInfo)
+      : Device(Dev), AsyncInfo(AsyncInfo) {}
+
+  /// A a private argument
+  int addArg(void *HstPtr, int64_t ArgSize, int64_t ArgOffset,
+             bool IsFirstPrivate, void *&TgtPtr, int TgtArgsIndex) {
+    // If the argument is not first-private, or its size is greater than a
+    // predefined threshold, we will allocate memory and issue the transfer
+    // immediately.
+    if (ArgSize > FirstPrivateArgSizeThreshold || !IsFirstPrivate) {
+      TgtPtr = Device.allocData(ArgSize, HstPtr);
+      if (!TgtPtr) {
+        DP("Data allocation for %sprivate array " DPxMOD " failed.\n",
+           (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtr));
+        return OFFLOAD_FAIL;
+      }
+#ifdef OMPTARGET_DEBUG
+      void *TgtPtrBase = (void *)((intptr_t)TgtPtr + ArgOffset);
+      DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD
+         " for %sprivate array " DPxMOD " - pushing target argument " DPxMOD
+         "\n",
+         ArgSize, DPxPTR(TgtPtr), (IsFirstPrivate ? "first-" : ""),
+         DPxPTR(HstPtr), DPxPTR(TgtPtrBase));
+#endif
+      // If first-private, copy data from host
+      if (IsFirstPrivate) {
+        int Ret = Device.submitData(TgtPtr, HstPtr, ArgSize, AsyncInfo);
+        if (Ret != OFFLOAD_SUCCESS) {
+          DP("Copying data to device failed, failed.\n");
+          return OFFLOAD_FAIL;
+        }
+      }
+      TgtPtrs.push_back(TgtPtr);
+    } else {
+      DP("Firstprivate array " DPxMOD " of size %" PRId64 " will be packed\n",
+         DPxPTR(HstPtr), ArgSize);
+      // When reach this point, the argument must meet all following
+      // requirements:
+      // 1. Its size does not exceed the threshold (see the comment for
+      // FirstPrivateArgSizeThreshold);
+      // 2. It must be first-private (needs to be mapped to target device).
+      // We will pack all this kind of arguments to transfer them all at once
+      // to reduce the number of data transfer. We will not take
+      // non-first-private arguments, aka. private arguments that doesn't need
+      // to be mapped to target device, into account because data allocation
+      // can be very efficient with memory manager.
+
+      // Placeholder value
+      TgtPtr = nullptr;
+      FirstPrivateArgInfo.emplace_back(TgtArgsIndex, HstPtr, ArgSize);
+      FirstPrivateArgSize += FirstPrivateArgInfo.back().AlignedSize;
+    }
+
+    return OFFLOAD_SUCCESS;
+  }
+
+  /// Pack first-private arguments, replace place holder pointers in \p TgtArgs,
+  /// and start the transfer.
+  int packAndTransfer(std::vector<void *> &TgtArgs) {
+    if (!FirstPrivateArgInfo.empty()) {
+      assert(FirstPrivateArgSize != 0 &&
+             "FirstPrivateArgSize is 0 but FirstPrivateArgInfo is empty");
+      FirstPrivateArgBuffer.resize(FirstPrivateArgSize, 0);
+      auto Itr = FirstPrivateArgBuffer.begin();
+      // Copy all host data to this buffer
+      for (FirstPrivateArgInfoTy &Info : FirstPrivateArgInfo) {
+        std::copy(Info.HstPtrBegin, Info.HstPtrEnd, Itr);
+        Itr = std::next(Itr, Info.AlignedSize);
+      }
+      // Allocate target memory
+      void *TgtPtr =
+          Device.allocData(FirstPrivateArgSize, FirstPrivateArgBuffer.data());
+      if (TgtPtr == nullptr) {
+        DP("Failed to allocate target memory for private arguments.\n");
+        return OFFLOAD_FAIL;
+      }
+      TgtPtrs.push_back(TgtPtr);
+      DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD "\n",
+         FirstPrivateArgSize, DPxPTR(TgtPtr));
+      // Transfer data to target device
+      int Ret = Device.submitData(TgtPtr, FirstPrivateArgBuffer.data(),
+                                  FirstPrivateArgSize, AsyncInfo);
+      if (Ret != OFFLOAD_SUCCESS) {
+        DP("Failed to submit data of private arguments.\n");
+        return OFFLOAD_FAIL;
+      }
+      // Fill in all placeholder pointers
+      auto TP = reinterpret_cast<uintptr_t>(TgtPtr);
+      for (FirstPrivateArgInfoTy &Info : FirstPrivateArgInfo) {
+        void *&Ptr = TgtArgs[Info.Index];
+        assert(Ptr == nullptr && "Target pointer is already set by mistaken");
+        Ptr = reinterpret_cast<void *>(TP);
+        TP += Info.AlignedSize;
+        DP("Firstprivate array " DPxMOD " of size %" PRId64 " mapped to " DPxMOD
+           "\n",
+           DPxPTR(Info.HstPtrBegin), Info.HstPtrEnd - Info.HstPtrBegin,
+           DPxPTR(Ptr));
+      }
+    }
+
+    return OFFLOAD_SUCCESS;
+  }
+
+  /// Free all target memory allocated for private arguments
+  int free() {
+    for (void *P : TgtPtrs) {
+      int Ret = Device.deleteData(P);
+      if (Ret != OFFLOAD_SUCCESS) {
+        DP("Deallocation of (first-)private arrays failed.\n");
+        return OFFLOAD_FAIL;
+      }
+    }
+
+    TgtPtrs.clear();
+
+    return OFFLOAD_SUCCESS;
+  }
+};
+
 /// Process data before launching the kernel, including calling targetDataBegin
 /// to map and transfer data to target device, transferring (first-)private
 /// variables.
 int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                       void **ArgBases, void **Args, int64_t *ArgSizes,
                       int64_t *ArgTypes, void **ArgMappers,
                       std::vector<void *> &TgtArgs,
                       std::vector<ptrdiff_t> &TgtOffsets,
-                      std::vector<void *> &FPArrays,
+                      PrivateArgumentManagerTy &PrivateArgumentManager,
                       __tgt_async_info *AsyncInfo) {
   DeviceTy &Device = Devices[DeviceId];
   int Ret = targetDataBegin(Device, ArgNum, ArgBases, Args, ArgSizes, ArgTypes,
                             ArgMappers, AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
     DP("Call to targetDataBegin failed, abort target.\n");
     return OFFLOAD_FAIL;
   }
for (int32_t I = 0; I < ArgNum; ++I) {
     ptrdiff_t TgtBaseOffset;
     bool IsLast, IsHostPtr; // unused.
     if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       DP("Forwarding first-private value " DPxMOD " to the target construct\n",
          DPxPTR(HstPtrBase));
       TgtPtrBegin = HstPtrBase;
       TgtBaseOffset = 0;
     } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {
-      // Allocate memory for (first-)private array
-      TgtPtrBegin = Device.allocData(ArgSizes[I], HstPtrBegin);
-      if (!TgtPtrBegin) {
-        DP("Data allocation for %sprivate array " DPxMOD " failed, "
-           "abort target.\n",
-           (ArgTypes[I] & OMP_TGT_MAPTYPE_TO ? "first-" : ""),
-           DPxPTR(HstPtrBegin));
-        return OFFLOAD_FAIL;
-      }
-      FPArrays.push_back(TgtPtrBegin);
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
-#ifdef OMPTARGET_DEBUG
-      void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
-      DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD " for "
-         "%sprivate array " DPxMOD " - pushing target argument " DPxMOD "\n",
-         ArgSizes[I], DPxPTR(TgtPtrBegin),
-         (ArgTypes[I] & OMP_TGT_MAPTYPE_TO ? "first-" : ""),
-         DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBase));
-#endif
-      // If first-private, copy data from host
-      if (ArgTypes[I] & OMP_TGT_MAPTYPE_TO) {
-        Ret =
-            Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSizes[I], AsyncInfo);
-        if (Ret != OFFLOAD_SUCCESS) {
-          DP("Copying data to device failed, failed.\n");
-          return OFFLOAD_FAIL;
-        }
-      }
+      const bool IsFirstPrivate = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
+      Ret = PrivateArgumentManager.addArg(HstPtrBegin, ArgSizes[I],
+                                          TgtBaseOffset, IsFirstPrivate,
+                                          TgtPtrBegin, TgtArgs.size());
+        DP("Failed to process %sprivate argument " DPxMOD "\n",
+           (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
     } else {
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
         HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
                                           false, IsHostPtr);
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
 #ifdef OMPTARGET_DEBUG
       void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
          DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));
 #endif
     }
     TgtArgsPositions[I] = TgtArgs.size();
     TgtArgs.push_back(TgtPtrBegin);
     TgtOffsets.push_back(TgtBaseOffset);
   }
 
   assert(TgtArgs.size() == TgtOffsets.size() &&
          "Size mismatch in arguments and offsets");
 
+  // Pack and transfer first-private arguments
+  Ret = PrivateArgumentManager.packAndTransfer(TgtArgs);
+  if (Ret != OFFLOAD_SUCCESS) {
+    DP("Failed to pack and transfer first private arguments\n");
+    return OFFLOAD_FAIL;
+  }
+
   return OFFLOAD_SUCCESS;
 }
 
 /// Process data after launching the kernel, including transferring data back to
 /// host if needed and deallocating target memory of (first-)private variables.
 int processDataAfter(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                      void **ArgBases, void **Args, int64_t *ArgSizes,
                      int64_t *ArgTypes, void **ArgMappers,
-                     std::vector<void *> &FPArrays,
+                     PrivateArgumentManagerTy &PrivateArgumentManager,
                      __tgt_async_info *AsyncInfo) {
   DeviceTy &Device = Devices[DeviceId];
 
   // Move data from device.
   int Ret = targetDataEnd(Device, ArgNum, ArgBases, Args, ArgSizes, ArgTypes,
                           ArgMappers, AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
     DP("Call to targetDataEnd failed, abort targe.\n");
     return OFFLOAD_FAIL;
   }
 
-  // Deallocate (first-)private arrays
-  for (void *P : FPArrays) {
-    Ret = Device.deleteData(P);
-    if (Ret != OFFLOAD_SUCCESS) {
-      DP("Deallocation of (first-)private arrays failed.\n");
-      return OFFLOAD_FAIL;
-    }
-  }
+  // Free target memory for private arguments
+  Ret = PrivateArgumentManager.free();
+    DP("Failed to deallocate target memory for private args\n");
 
   return OFFLOAD_SUCCESS;
 }
 } // namespace
 
 /// performs the same actions as data_begin in case arg_num is
 /// non-zero and initiates run of the offloaded region on the target platform;
 /// if arg_num is non-zero after the region execution is done it also
__tgt_target_table *TargetTable = nullptr;
     TargetTable = TM->Table->TargetsTable[DeviceId];
   }
   assert(TargetTable && "Global data has not been mapped\n");
 
   __tgt_async_info AsyncInfo;
 
   std::vector<void *> TgtArgs;
   std::vector<ptrdiff_t> TgtOffsets;
-  std::vector<void *> FPArrays;
+
+  PrivateArgumentManagerTy PrivateArgumentManager(Device, &AsyncInfo);
 
   // Process data, such as data mapping, before launching the kernel
   int Ret = processDataBefore(DeviceId, HostPtr, ArgNum, ArgBases, Args,
                               ArgSizes, ArgTypes, ArgMappers, TgtArgs,
-                              TgtOffsets, FPArrays, &AsyncInfo);
+                              TgtOffsets, PrivateArgumentManager, &AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
     DP("Failed to process data before launching the kernel.\n");
     return OFFLOAD_FAIL;
   }
 
   // Get loop trip count
   uint64_t LoopTripCount = getLoopTripCount(DeviceId);
 
int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
   if (Ret != OFFLOAD_SUCCESS) {
     DP("Executing target region abort target.\n");
     return OFFLOAD_FAIL;
   }
 
   // Transfer data back and deallocate target memory for (first-)private
   // variables
   Ret = processDataAfter(DeviceId, HostPtr, ArgNum, ArgBases, Args, ArgSizes,
-                         ArgTypes, ArgMappers, FPArrays, &AsyncInfo);
+                         ArgTypes, ArgMappers, PrivateArgumentManager,
+                         &AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
     DP("Failed to process data after launching the kernel.\n");
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }

openmp/libomptarget/test/mapping/private_mapping.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda
+// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 \
+// RUN: | %fcheck-nvptx64-nvidia-cuda
+
+#include <assert.h>
+#include <stdio.h>
+
+int main() {
+  int data1[3] = {1}, data2[3] = {2}, data3[3] = {3};
+  int sum[16] = {0};
+#pragma omp target teams distribute parallel for map(tofrom                    \
+                                                     : sum)                    \
+    firstprivate(data1, data2, data3)
+  for (int i = 0; i < 16; ++i) {
+    for (int j = 0; j < 3; ++j) {
+      sum[i] += data1[j];
+      sum[i] += data2[j];
+      sum[i] += data3[j];
+    }
+  }
+
+  for (int i = 0; i < 16; ++i) {
+    assert(sum[i] == 6);
+  }
+
+  printf("PASS\n");
+
+  return 0;
+}
+
+// CHECK: PASS