From: Thomas Schwinge <tschwinge@baylibre.com>
Date: Mon, 3 Nov 2025 14:24:28 +0000 (+0100)
Subject: Fix 'libgomp.c/pr122281.c' for non-USM offloading execution [PR122281]
X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=7a25badc5ddd95df53a3ca5156d9dd071f3f78c1;p=thirdparty%2Fgcc.git

Fix 'libgomp.c/pr122281.c' for non-USM offloading execution [PR122281]

... where it currently runs into:

    libgomp: cuCtxSynchronize error: an illegal memory access was encountered

... for nvptx, or similarly for GCN:

    Memory access fault by GPU node-1 (Agent handle: 0x34d77290) on address 0x7fff3c553000. Reason: Page not present or supervisor privilege.

Fix-up for commit r16-4961-ge2cbcd1b27c0da92bdcd96664064d3d0c1d44e6f
"Fix gimple_copy for OpenMP atomic load/store [PR122281, PR105001]".

	PR libgomp/122281
	libgomp/
	* testsuite/libgomp.c/pr122281.c: Fix for non-USM offloading
	execution.
---

diff --git a/libgomp/testsuite/libgomp.c/pr122281.c b/libgomp/testsuite/libgomp.c/pr122281.c
index a02a728c5f0..68fc3be7221 100644
--- a/libgomp/testsuite/libgomp.c/pr122281.c
+++ b/libgomp/testsuite/libgomp.c/pr122281.c
@@ -18,7 +18,7 @@
 int __attribute__((noinline, noclone))
 f(int *A, int n, int dev) {
  int cnt = 0;
- #pragma omp target map(cnt) device(dev)
+ #pragma omp target map(cnt) map(to:A[0:n]) device(dev)
  {
    #pragma omp parallel for simd
    for (int i = 0; i < n; i++)