diff options
author | Johannes Doerfert <johannes@jdoerfert.de> | 2023-03-21 15:48:11 -0700 |
---|---|---|
committer | Johannes Doerfert <johannes@jdoerfert.de> | 2023-03-21 19:16:13 -0700 |
commit | de9edf4afecc1c2caf3b552f9241008ad2bd40a8 (patch) | |
tree | 0e4eec083b1082791fe710c372d31dc72463b390 /openmp | |
parent | ebcc6dba5f0815877322256095b400b31adac5f4 (diff) | |
download | llvm-de9edf4afecc1c2caf3b552f9241008ad2bd40a8.tar.gz |
[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
Diffstat (limited to 'openmp')
-rw-r--r-- | openmp/libomptarget/src/device.cpp | 3 | ||||
-rw-r--r-- | openmp/libomptarget/src/omptarget.cpp | 2 | ||||
-rw-r--r-- | openmp/libomptarget/test/mapping/data_member_ref.cpp | 34 |
3 files changed, 25 insertions, 14 deletions
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index d670bad1342b..1f5d5a23371d 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 392a9f79bd76..2158b948bc9e 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 6b52a04e34f1..fdb8abcaa650 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; |