diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 7cf411155c6..42591e4029a 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -14258,6 +14258,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      if (OMP_CLAUSE_MAP_READONLY (c))
+	OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
       OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	  && !c_mark_addressable (t))
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 8fb47fd179e..6ab467e1140 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5872,6 +5872,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	    }
 	  else
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  if (OMP_CLAUSE_MAP_READONLY (c))
+	    OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
 	  OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
 	  if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	      && !cxx_mark_addressable (t))
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 2253d559f9c..d7cd65af1bb 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2524,6 +2524,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
       node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
       OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
+      if (n->u.readonly)
+	OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
       /* This purposely does not include GOMP_MAP_ALWAYS_POINTER.  The extra
 	 cast prevents gimplify.cc from recognising it as being part of the
 	 struct - and adding an 'alloc: for the 'desc.data' pointer, which
@@ -2559,6 +2561,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
 				OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
       OMP_CLAUSE_DECL (node3) = decl;
+      if (n->u.readonly)
+	OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
     }
   ptr2 = fold_convert (ptrdiff_type_node, ptr2);
   OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
diff --git a/gcc/gimple-expr.cc b/gcc/gimple-expr.cc
index f15cc0ba715..42c0f6469b1 100644
--- a/gcc/gimple-expr.cc
+++ b/gcc/gimple-expr.cc
@@ -376,6 +376,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))
+    DECL_POINTS_TO_READONLY (copy) = DECL_POINTS_TO_READONLY (var);
   DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var);
   if (DECL_USER_ALIGN (var))
     {
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 36e5df050b9..394e40fead2 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -221,6 +221,7 @@ struct gimplify_omp_ctx
   splay_tree variables;
   hash_set<tree> *privatized_types;
   tree clauses;
+  hash_set<tree_operand_hash> *pt_readonly_ptrs;
   /* Iteration variables in an OMP_FOR.  */
   vec<tree> loop_iter_var;
   location_t location;
@@ -628,6 +629,15 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
   gimplify_expr (&val, pre_p, post_p, is_gimple_reg_rhs_or_call,
 		 fb_rvalue);
 
+  bool pt_readonly = false;
+  if (gimplify_omp_ctxp && gimplify_omp_ctxp->pt_readonly_ptrs)
+    {
+      tree ptr = val;
+      if (TREE_CODE (ptr) == POINTER_PLUS_EXPR)
+	ptr = TREE_OPERAND (ptr, 0);
+      pt_readonly = gimplify_omp_ctxp->pt_readonly_ptrs->contains (ptr);
+    }
+
   if (allow_ssa
       && gimplify_ctxp->into_ssa
       && is_gimple_reg_type (TREE_TYPE (val)))
@@ -639,9 +649,18 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
 	  if (name)
 	    SET_SSA_NAME_VAR_OR_IDENTIFIER (t, create_tmp_var_name (name));
 	}
+      if (pt_readonly)
+	SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
     }
   else
-    t = lookup_tmp_var (val, is_formal, not_gimple_reg);
+    {
+      t = lookup_tmp_var (val, is_formal, not_gimple_reg);
+      if (pt_readonly)
+	{
+	  DECL_POINTS_TO_READONLY (t) = 1;
+	  gimplify_omp_ctxp->pt_readonly_ptrs->add (t);
+	}
+    }
 
   mod = build2 (INIT_EXPR, TREE_TYPE (t), t, unshare_expr (val));
 
@@ -8906,6 +8925,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
   OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
   OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end));
   OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
+  if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (grp_end))
+    OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
   tree grp_mid = NULL_TREE;
   if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
     grp_mid = OMP_CLAUSE_CHAIN (grp_start);
@@ -11741,6 +11762,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	      gimplify_omp_ctxp = outer_ctx;
 	    }
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		   && (code == OACC_PARALLEL
+		       || code == OACC_KERNELS
+		       || code == OACC_SERIAL)
+		   && OMP_CLAUSE_MAP_POINTS_TO_READONLY (c))
+	    {
+	      if (ctx->pt_readonly_ptrs == NULL)
+		ctx->pt_readonly_ptrs = new hash_set<tree_operand_hash> ();
+	      ctx->pt_readonly_ptrs->add (OMP_CLAUSE_DECL (c));
+	    }
 	  if (notice_outer)
 	    goto do_notice;
 	  break;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index b882df048ef..204fc72ca2d 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -14098,6 +14098,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))
+		  DECL_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 171f96c08db..1f10fd25e46 100644
--- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -19,8 +19,8 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "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 00000000000..d32d3362000
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
@@ -0,0 +1,15 @@
+/* { dg-additional-options "-O -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\\\];" 1 "fre1" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
index 069fec0a0d5..1e5e60f9744 100644
--- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -20,8 +20,8 @@ program main
   !$acc end parallel
 end program main
 
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) .+ map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) .+ map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a.0 \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) a.0\\\]\\) map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) b\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\\]\\) map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\\]\\)" 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: .+\\\]\\);" 2 "original" } }
 
 
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 9604c3eecc5..1a8b121f30b 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -907,6 +907,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:
@@ -3436,6 +3438,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 23387b90fe3..32d35a29dfc 100644
--- a/gcc/tree-ssanames.cc
+++ b/gcc/tree-ssanames.cc
@@ -402,6 +402,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple *stmt,
   else
     SSA_NAME_RANGE_INFO (t) = NULL;
 
+  if (VAR_P (var) && DECL_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 ac563de1fc3..880ffb367a3 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1021,6 +1021,13 @@ extern void omp_clause_range_check_failed (const_tree, const char *, int,
 #define DECL_HIDDEN_STRING_LENGTH(NODE) \
   (TREE_CHECK (NODE, PARM_DECL)->decl_common.decl_nonshareable_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 DECL_POINTS_TO_READONLY(NODE) \
+  (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
+
 /* In a CALL_EXPR, means that the call is the jump from a thunk to the
    thunked-to function.  Be careful to avoid using this macro when one of the
    next two applies instead.  */
@@ -1815,6 +1822,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_))
