On 2024/5/16 8:36 PM, Richard Biener wrote:
>> After omp-expand (before SSA):
>>
>> __attribute__((oacc parallel, omp target entrypoint, noclone))
>> void main._omp_fn.1 (const struct .omp_data_t.3 & restrict .omp_data_i)
>> {
>> ...
>> <bb 5> :
>> D.2962 = .omp_data_i->D.2947;
>> a.8 = D.2962;
> So 'readonly: a[:32]' is put in .omp_data_i->D.2947 in the caller
> and extracted here. And you arrange for 'a.8' to have
> DECL_POINTS_TO_READONLY set by "magic"? Looking at this I wonder
> if it would be more useful to "const qualify" (but "really", not
> in the C sense) .omp_data_i->D.2947 instead? Thus have a
> FIELD_POINTS_TO_READONLY_MEMORY flag on the FIELD_DECL.
>
> Points-to analysis should then be able to handle this similar to how
> it handles loads of restrict qualified pointers. Well, of course not
> as simple since it now adds "qualifiers" to storage since I presume
> the same object can be both readonly and not readonly like via
>
> #pragma acc parallel copyin(readonly: a[:32], a[33:64]) copyout(r)
>
> ? That is, currently there's only one "readonly" object kind in
> points-to, that's STRING_CSTs which get all globbed to string_id
> and "ignored" for alias purposes since you can't change them.
>
> So possibly you want to combine this with restrict qualifying the
> pointer so we know there's no other (read-write) access to the memory
> possible. But then you might get all the good stuff already by
> _just_ doing that restrict qualification and ignoring the readonly-ness?
Hi Richard,
(sorry for returning to this with such a long hiatus)
The restrict qualification on pointers is interesting, I'll investigate using
that
in a later revision of this patch.
> +#define DECL_POINTS_TO_READONLY(NODE) \
> + (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
>
> you need to document uses of flags in tree-core.h to avoid clashes.
> Also since this doesn't apply to all DECLs it should be named
> VAR_POINTS_TO_...
I've made that change, seems logical. Though frankly speaking there appears to
be
many DECL_* macros that have "TREE_CHECK (NODE, VAR_DECL)" :P
> I still think this is too fragile - there's no real constraints
> on what VAR_DECL we create SSA names off, so the automatism
> in make_ssa_name_fn and esp. copy_var_decl and via copy_node
> copy_decl_no_change, thus during inlining, makes your arguments
> only apply to the use for OpenMP - but nothing above hints at
> this is just usable there, asking for trouble.
The idea was that we restrict what VAR_DECLs we allow such a points-to-readonly
set,
but yeah this could use more documentation later.
I have for the time being pushed the attached patch to devel/omp/gcc-14 branch.
Thanks,
Chung-Lin
From 87dcdc3deb6ab6b1bc09aacbfb828182afe79f74 Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <[email protected]>
Date: Mon, 11 Nov 2024 17:16:26 +0000
Subject: [PATCH] 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.
---
gcc/c-family/c-omp.cc | 4 ++++
gcc/fortran/trans-openmp.cc | 2 ++
gcc/gimple-expr.cc | 2 ++
gcc/omp-low.cc | 2 ++
gcc/testsuite/c-c++-common/goacc/readonly-1.c | 20 +++++++++----------
gcc/testsuite/c-c++-common/goacc/readonly-2.c | 16 +++++++++++++++
.../gfortran.dg/goacc/readonly-1.f90 | 20 +++++++++----------
gcc/testsuite/gfortran.dg/pr67170.f90 | 2 +-
gcc/tree-pretty-print.cc | 4 ++++
gcc/tree-ssanames.cc | 3 +++
gcc/tree.h | 11 ++++++++++
11 files changed, 65 insertions(+), 21 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/readonly-2.c
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 6b10082cc93..479b69741a5 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4135,6 +4135,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)
@@ -4324,6 +4326,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 ff7a8d19c00..e6ba7b5839e 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4032,6 +4032,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 f8d7185530c..22722249a19 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 a190e0fb2eb..fa24feb80bb 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -15293,6 +15293,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 bd4d2294c2a..aeb8e0ef862 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 00000000000..3f52a9f6afb
--- /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 fc1e2719e67..cad449e6d40 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 80236470f42..d7c33a4c3db 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 2a98cb0986b..bc4847a6b8f 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -929,6 +929,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:
@@ -3738,6 +3740,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 6c2525900ab..2ac2515e9c8 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) && 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 48c548ab585..53b1aed4435 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1867,6 +1867,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_))
@@ -3347,6 +3351,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)
--
2.34.1
