OpenACC 2.7: Connect readonly modifier to points-to analysis

This patch links the readonly modifier to points-to analysis.

In front-ends, firstprivate pointer clauses are marked with
OMP_CLAUSE_MAP_POINTS_TO_READONLY set true, and later during lowering the
receiver side read of pointer has VAR_POINTS_TO_READONLY set true, which later
directs SSA_NAME_POINTS_TO_READONLY_MEMORY set to true during SSA conversion.

SSA_NAME_POINTS_TO_READONLY_MEMORY is an already existing flag connected with
alias oracle routines in tree-ssa-alias.cc, thus making the readonly-modifier
effective in hinting points-to analysis.

Currently have one testcase c-c++-common/goacc/readonly-2.c where we can
demonstrate 'readonly' can avoid a clobber by function call.

This patch is ported from upstream submission:
https://gcc.gnu.org/pipermail/gcc-patches/2024-April/648728.html

gcc/c-family/ChangeLog:

	* c-omp.cc (c_omp_address_inspector::expand_array_base):
	Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
	(c_omp_address_inspector::expand_component_selector): Likewise.

gcc/fortran/ChangeLog:

	* trans-openmp.cc (gfc_trans_omp_array_section):
	Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.

gcc/ChangeLog:

	* gimple-expr.cc (copy_var_decl): Copy VAR_POINTS_TO_READONLY
	for VAR_DECLs.
	* omp-low.cc (lower_omp_target): Set VAR_POINTS_TO_READONLY for
	variables of receiver refs.
	* tree-pretty-print.cc (dump_omp_clause):
	Print OMP_CLAUSE_MAP_POINTS_TO_READONLY.
	(dump_generic_node): Print SSA_NAME_POINTS_TO_READONLY_MEMORY.
	* tree-ssanames.cc (make_ssa_name_fn): Set
	SSA_NAME_POINTS_TO_READONLY_MEMORY if DECL_POINTS_TO_READONLY is set.
	* tree.h (OMP_CLAUSE_MAP_POINTS_TO_READONLY): New macro.
	(VAR_POINTS_TO_READONLY): New macro.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/readonly-1.c: Adjust testcase.
	* c-c++-common/goacc/readonly-2.c: New testcase.
	* gfortran.dg/goacc/readonly-1.f90: Adjust testcase.
	* gfortran.dg/pr67170.f90: Likewise.

11 files changed, 65 insertions, 21 deletions

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 588218d..662b791 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4157,6 +4157,8 @@ c_omp_address_inspector::expand_array_base (tree *pc,
     }
   else if (c2)
     {
+      if (OMP_CLAUSE_MAP_READONLY (c))
+	OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
       OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
       OMP_CLAUSE_CHAIN (c) = c2;
       if (implicit_p)
@@ -4346,6 +4348,8 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
     }
   else if (c2)
     {
+      if (OMP_CLAUSE_MAP_READONLY (c))
+	OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
       OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
       OMP_CLAUSE_CHAIN (c) = c2;
       pc = &OMP_CLAUSE_CHAIN (c);
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index afe18ec..521f7cc 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3408,6 +3408,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, toc_directive cd,
   ptr2 = fold_convert (ptrdiff_type_node, ptr2);
   OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
 					 ptr, ptr2);
+  if (n->u.map.readonly)
+    OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
 }
 
 /* CLAUSES is a list of clauses resulting from an "omp declare mapper"
diff --git a/gcc/gimple-expr.cc b/gcc/gimple-expr.cc
index c0367f4..69a5b7e 100644
--- a/gcc/gimple-expr.cc
+++ b/gcc/gimple-expr.cc
@@ -385,6 +385,8 @@ copy_var_decl (tree var, tree name, tree type)
   DECL_CONTEXT (copy) = DECL_CONTEXT (var);
   TREE_USED (copy) = 1;
   DECL_SEEN_IN_BIND_EXPR_P (copy) = 1;
+  if (VAR_P (var))
+    VAR_POINTS_TO_READONLY (copy) = VAR_POINTS_TO_READONLY (var);
   DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var);
   if (DECL_USER_ALIGN (var))
     {
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 7689d62..ee2b1f3 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -15267,6 +15267,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		if (ref_to_array)
 		  x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
 		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (c) && VAR_P (x))
+		  VAR_POINTS_TO_READONLY (x) = 1;
 		if ((is_ref && !ref_to_array)
 		    || ref_to_ptr)
 		  {
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
index bd4d229..aeb8e0e 100644
--- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -48,17 +48,17 @@ int main (void)
 
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc declare map\\(to:y\\) map\\(readonly,to:s\\) map\\(readonly,to:x\\)" 1 "original" } } */
 
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
 
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-2.c b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
new file mode 100644
index 0000000..3f52a9f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
@@ -0,0 +1,16 @@
+/* { dg-additional-options "-O -fdump-tree-phiprop -fdump-tree-fre" } */
+
+#pragma acc routine
+extern void foo (int *ptr, int val);
+
+int main (void)
+{
+  int r, a[32];
+  #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
+  {
+    foo (a, a[8]);
+    r = a[8];
+  }
+}
+/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ = MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 2 "phiprop1" } } */
+/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ = MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 1 "fre1" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
index fc1e271..cad449e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -80,16 +80,16 @@ end program main
 ! The front end turns OpenACC 'declare' into OpenACC 'data'.
 !   { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*b\\) map\\(alloc:b.+ map\\(to:\\*c\\) map\\(alloc:c.+" 1 "original" } }
 !   { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:g\\) map\\(to:h\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:b.+ map\\(pt_readonly,alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:b.+ map\\(pt_readonly,alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:b.+ map\\(pt_readonly,alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:b.+ map\\(pt_readonly,alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:\\*.+ map\\(pt_readonly,alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(pt_readonly,alloc:a.+ map\\(readonly,to:b.+ map\\(pt_readonly,alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
 
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/pr67170.f90 b/gcc/testsuite/gfortran.dg/pr67170.f90
index 8023647..d7c33a4 100644
--- a/gcc/testsuite/gfortran.dg/pr67170.f90
+++ b/gcc/testsuite/gfortran.dg/pr67170.f90
@@ -28,4 +28,4 @@ end subroutine foo
 end module test_module
 end program
 
-! { dg-final { scan-tree-dump-times "= \\*arg_\[0-9\]+\\(D\\);" 1 "fre1" } }
+! { dg-final { scan-tree-dump-times "= \\*arg_\[0-9\]+\\(D\\)\\(ptro\\);" 1 "fre1" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 2826da7..6140d51 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1026,6 +1026,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       pp_string (pp, "map(");
       if (OMP_CLAUSE_MAP_READONLY (clause))
 	pp_string (pp, "readonly,");
+      if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause))
+	pp_string (pp, "pt_readonly,");
       switch (OMP_CLAUSE_MAP_KIND (clause))
 	{
 	case GOMP_MAP_ALLOC:
@@ -3861,6 +3863,8 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
 	pp_string (pp, "(D)");
       if (SSA_NAME_OCCURS_IN_ABNORMAL_PHI (node))
 	pp_string (pp, "(ab)");
+      if (SSA_NAME_POINTS_TO_READONLY_MEMORY (node))
+	pp_string (pp, "(ptro)");
       break;
 
     case WITH_SIZE_EXPR:
diff --git a/gcc/tree-ssanames.cc b/gcc/tree-ssanames.cc
index d7865f2..3605c4d 100644
--- a/gcc/tree-ssanames.cc
+++ b/gcc/tree-ssanames.cc
@@ -403,6 +403,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple *stmt,
   else
     SSA_NAME_RANGE_INFO (t) = NULL;
 
+  if (VAR_P (var) && VAR_POINTS_TO_READONLY (var))
+    SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
+
   SSA_NAME_IN_FREE_LIST (t) = 0;
   SSA_NAME_IS_DEFAULT_DEF (t) = 0;
   init_ssa_name_imm_use (t);
diff --git a/gcc/tree.h b/gcc/tree.h
index 40d8930..28b6089 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1905,6 +1905,10 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_MAP_READONLY(NODE) \
   TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
 
+/* Set if 'OMP_CLAUSE_DECL (NODE)' points to read-only memory.  */
+#define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \
+  TREE_CONSTANT (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
 /* Same as above, for use in OpenACC cache directives.  */
 #define OMP_CLAUSE__CACHE__READONLY(NODE) \
   TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_))
@@ -3398,6 +3402,13 @@ extern void decl_fini_priority_insert (tree, priority_type);
 #define VAR_DECL_IS_VIRTUAL_OPERAND(NODE) \
   (VAR_DECL_CHECK (NODE)->base.u.bits.saturating_flag)
 
+/* In a VAR_DECL, set for variables regarded as pointing to memory not written
+   to. SSA_NAME_POINTS_TO_READONLY_MEMORY gets set for SSA_NAMEs created from
+   such VAR_DECLs. Currently used by OpenACC 'readonly' modifier in copyin
+   clauses.  */
+#define VAR_POINTS_TO_READONLY(NODE) \
+  (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
+
 /* In a VAR_DECL, nonzero if this is a non-local frame structure.  */
 #define DECL_NONLOCAL_FRAME(NODE)  \
   (VAR_DECL_CHECK (NODE)->base.default_def_flag)