[Openmp-commits] [openmp] de9edf4 - [OpenMP] Avoid zero size copies to the device

[Johannes Doerfert via Openmp-commits]

Author: Johannes Doerfert
Date: 2023-03-21T19:16:13-07:00
New Revision: de9edf4afecc1c2caf3b552f9241008ad2bd40a8

URL: https://github.com/llvm/llvm-project/commit/de9edf4afecc1c2caf3b552f9241008ad2bd40a8
DIFF: https://github.com/llvm/llvm-project/commit/de9edf4afecc1c2caf3b552f9241008ad2bd40a8.diff

LOG: [OpenMP] Avoid zero size copies to the device

This unblocks one of the XFAIL tests for AMD, though we need to work
around the missing printf still.

Differential Revision: https://reviews.llvm.org/D146592

Added: 
    

Modified: 
    openmp/libomptarget/src/device.cpp
    openmp/libomptarget/src/omptarget.cpp
    openmp/libomptarget/test/mapping/data_member_ref.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index d670bad1342bc..1f5d5a23371db 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -321,7 +321,8 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
 
   // If the target pointer is valid, and we need to transfer data, issue the
   // data transfer.
-  if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {
+  if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways) &&
+      Size != 0) {
     // Lock the entry before releasing the mapping table lock such that another
     // thread that could issue data movement will get the right result.
     std::lock_guard<decltype(*Entry)> LG(*Entry);

diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 392a9f79bd760..2158b948bc9ec 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -1028,7 +1028,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     // Move data back to the host
     const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
     const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
-    if (HasFrom && (HasAlways || IsLast) && !IsHostPtr) {
+    if (HasFrom && (HasAlways || IsLast) && !IsHostPtr && DataSize != 0) {
       DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
          DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
 

diff  --git a/openmp/libomptarget/test/mapping/data_member_ref.cpp b/openmp/libomptarget/test/mapping/data_member_ref.cpp
index 6b52a04e34f1d..fdb8abcaa6506 100644
--- a/openmp/libomptarget/test/mapping/data_member_ref.cpp
+++ b/openmp/libomptarget/test/mapping/data_member_ref.cpp
@@ -1,8 +1,5 @@
 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <stdio.h>
 
 struct View {
@@ -26,42 +23,55 @@ int main() {
   int Data = 123456;
   V1.Data = &Data;
   Foo<ViewPtr> Baz(V1);
+  int D1, D2;
 
   // CHECK: Host 123456.
   printf("Host %d.\n", Bar.VRef.Data);
-#pragma omp target map(Bar.VRef)
+#pragma omp target map(Bar.VRef) map(from : D1, D2)
   {
     // CHECK: Device 123456.
-    printf("Device %d.\n", Bar.VRef.Data);
+    D1 = Bar.VRef.Data;
+    printf("Device %d.\n", D1);
     V.Data = 654321;
     // CHECK: Device 654321.
-    printf("Device %d.\n", Bar.VRef.Data);
+    D2 = Bar.VRef.Data;
+    printf("Device %d.\n", D2);
   }
+  printf("Device %d.\n", D1);
+  printf("Device %d.\n", D2);
   // CHECK: Host 654321 654321.
   printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
   V.Data = 123456;
   // CHECK: Host 123456.
   printf("Host %d.\n", Bar.VRef.Data);
-#pragma omp target map(Bar) map(Bar.VRef)
+#pragma omp target map(Bar) map(Bar.VRef) map(from : D1, D2)
   {
     // CHECK: Device 123456.
-    printf("Device %d.\n", Bar.VRef.Data);
+    D1 = Bar.VRef.Data;
+    printf("Device %d.\n", D1);
     V.Data = 654321;
     // CHECK: Device 654321.
-    printf("Device %d.\n", Bar.VRef.Data);
+    D2 = Bar.VRef.Data;
+    printf("Device %d.\n", D2);
   }
+  printf("Device %d.\n", D1);
+  printf("Device %d.\n", D2);
   // CHECK: Host 654321 654321.
   printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
   // CHECK: Host 123456.
   printf("Host %d.\n", *Baz.VRef.Data);
-#pragma omp target map(*Baz.VRef.Data)
+#pragma omp target map(*Baz.VRef.Data) map(from : D1, D2)
   {
     // CHECK: Device 123456.
-    printf("Device %d.\n", *Baz.VRef.Data);
+    D1 = *Baz.VRef.Data;
+    printf("Device %d.\n", D1);
     *V1.Data = 654321;
     // CHECK: Device 654321.
-    printf("Device %d.\n", *Baz.VRef.Data);
+    D2 = *Baz.VRef.Data;
+    printf("Device %d.\n", D2);
   }
+  printf("Device %d.\n", D1);
+  printf("Device %d.\n", D2);
   // CHECK: Host 654321 654321 654321.
   printf("Host %d %d %d.\n", *Baz.VRef.Data, *V1.Data, Data);
   return 0;