https://gcc.gnu.org/g:58925a5a130561aa89fd20ea4f731f2c73384b2e
commit r17-2444-g58925a5a130561aa89fd20ea4f731f2c73384b2e Author: Thomas Schwinge <[email protected]> Date: Thu Jul 16 00:47:08 2026 +0200 OpenMP: Handle 'device_type(nohost)' on 'target': Fix 'libgomp.c/target-device-type-1.c' for non-USM 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. Diff: --- 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; }
