With commit 97a43fc4dbc5cfb1174a13313dda7637c838d110
"OpenMP: Handle 'device_type(nohost)' on 'target'", I'm seeing:
PASS: libgomp.c/target-device-type-1.c (test for excess errors)
FAIL: libgomp.c/target-device-type-1.c execution test
PASS: libgomp.c/target-device-type-1.c scan-tree-dump gimple "=
__builtin_omp_is_initial_device \\(\\);"
PASS: libgomp.c/target-device-type-1.c scan-tree-dump gimple
"__builtin_GOMP_error \\("Executing device-type ..nohost.. target region on the
host","
... for (non-USM) AMD GPU offloading:
Memory access fault by GPU node-2 (Agent handle: 0xff37f00) on address
0x7ffd8a413000. Reason: Page not present or supervisor privilege.
..., and correspondingly for (non-USM) NVIDIA GPU offloading:
libgomp: cuCtxSynchronize error: an illegal memory access was encountered
Fix this by 'map'ping 'x' for the OpenMP 'target' construct. Not technically
necessary for 'libgomp.c/target-device-type-2.c' (as that one's exercising the
'host' error case), but still apply the same change also to that one.
libgomp/
* testsuite/libgomp.c/target-device-type-1.c (f): Map 'x'.
* testsuite/libgomp.c/target-device-type-2.c (f): Likewise.
---
libgomp/testsuite/libgomp.c/target-device-type-1.c | 2 +-
libgomp/testsuite/libgomp.c/target-device-type-2.c | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/libgomp/testsuite/libgomp.c/target-device-type-1.c
b/libgomp/testsuite/libgomp.c/target-device-type-1.c
index a1e1af6979ad..d8d964e7fe5c 100644
--- a/libgomp/testsuite/libgomp.c/target-device-type-1.c
+++ b/libgomp/testsuite/libgomp.c/target-device-type-1.c
@@ -4,7 +4,7 @@
int
f (int *x)
{
- #pragma omp target device_type(nohost)
+ #pragma omp target map(tofrom:x[0:1]) device_type(nohost)
(*x)++;
return *x;
}
diff --git a/libgomp/testsuite/libgomp.c/target-device-type-2.c
b/libgomp/testsuite/libgomp.c/target-device-type-2.c
index 37fa62216c93..54550422032c 100644
--- a/libgomp/testsuite/libgomp.c/target-device-type-2.c
+++ b/libgomp/testsuite/libgomp.c/target-device-type-2.c
@@ -6,7 +6,7 @@
int
f (int *x)
{
- #pragma omp target device_type(nohost)
+ #pragma omp target map(tofrom:x[0:1]) device_type(nohost)
(*x)++;
return *x;
}
--
2.34.1
