summaryrefslogtreecommitdiff
path: root/openmp
diff options
context:
space:
mode:
authorJohannes Doerfert <johannes@jdoerfert.de>2023-03-21 15:48:11 -0700
committerJohannes Doerfert <johannes@jdoerfert.de>2023-03-21 19:16:13 -0700
commitde9edf4afecc1c2caf3b552f9241008ad2bd40a8 (patch)
tree0e4eec083b1082791fe710c372d31dc72463b390 /openmp
parentebcc6dba5f0815877322256095b400b31adac5f4 (diff)
downloadllvm-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.cpp3
-rw-r--r--openmp/libomptarget/src/omptarget.cpp2
-rw-r--r--openmp/libomptarget/test/mapping/data_member_ref.cpp34
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;