[PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives

Julian Brown julian@codesourcery.com
Tue Jun 16 22:39:43 GMT 2020


When attaching pointers in Fortran, OpenACC 2.6 specifies that a
descriptor must be copied to the target at the same time (see next
patch).  That means that stripping GOMP_MAP_TO_PSET (and lesserly,
GOMP_MAP_POINTER), which was behaviour introduced by the manual deep-copy
middle-end support patch, was probably wrong.

That arguably answers some of the questions at the end of:

https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547424.html

It appears that the user can (but certainly should not!) map a synthesized
array descriptor using an "enter data" operation that can go out of
scope before that data is unmapped.  It would be nice to give a warning
for an attempt to do such a thing, though I have no idea if that's
possible in practice.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): Do not strip
	GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data
	directives.

	gcc/testsuite/
	* gfortran.dg/goacc/finalize-1.f: Update expected dump output.
---
 gcc/gimplify.c                               | 11 ++---------
 gcc/testsuite/gfortran.dg/goacc/finalize-1.f |  4 ++--
 2 files changed, 4 insertions(+), 11 deletions(-)

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9851edfc4db..aa6853f0dcc 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8767,6 +8767,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
+	    case OACC_ENTER_DATA:
+	    case OACC_EXIT_DATA:
 	    case OACC_HOST_DATA:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
@@ -8775,15 +8777,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		   mapped, but not the pointer to it.  */
 		remove = true;
 	      break;
-	    case OACC_ENTER_DATA:
-	    case OACC_EXIT_DATA:
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
-		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
-		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
-		  || (OMP_CLAUSE_MAP_KIND (c)
-		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
-		remove = true;
-	      break;
 	    default:
 	      break;
 	    }
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 1e2e3e94b8a..ca642156e9f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -21,7 +21,7 @@
 
 !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -33,5 +33,5 @@
 
 !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
-- 
2.23.0



More information about the Gcc-patches mailing list