On 2023/7/11 2:33 AM, Chung-Lin Tang via Gcc-patches wrote:
> As we discussed earlier, the work for actually linking this to middle-end
> points-to analysis is a somewhat non-trivial issue. This first patch allows
> the language feature to be used in OpenACC directives first (with no effect
> for now).
> The middle-end changes are probably going to be a later patch.
This second patch tries to link the readonly modifier to points-to analysis.
There already exists SSA_NAME_POINTS_TO_READONLY_MEMORY and it's support in the
alias oracle routines in tree-ssa-alias.cc, so basically what this patch does is
try to make the variables holding the array section base pointers to have this
flag set.
There is an another OMP_CLAUSE_MAP_POINTS_TO_READONLY set by front-ends on the
associated pointer clauses if OMP_CLAUSE_MAP_READONLY is set.
Also a DECL_POINTS_TO_READONLY flag is set for VAR_DECLs when creating the tmp
vars carrying these receiver references on the offloaded side. These
eventually get translated to SSA_NAME_POINTS_TO_READONLY_MEMORY.
This still doesn't always work as expected in terms of optimization:
struct pointer fields and Fortran arrays (kind of like C structs) which have
several accesses to create the pointer access on the receive/offloaded side,
and SRA appears to not work on these sequences, so gets in the way of much
redundancy elimination.
Currently have one testcase where we can demonstrate 'readonly' can avoid
a clobber by function call. Tested on powerpc64le-linux/nvptx.
Note this patch is create a-top of the front-end patch.
(will respond to the other front-end patch comments later)
Thanks,
Chung-Lin
2023-07-25 Chung-Lin Tang <clt...@codesourcery.com>
gcc/c/ChangeLog:
* c-typeck.cc (handle_omp_array_sections):
Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
gcc/cp/ChangeLog:
* semantics.cc (handle_omp_array_sections):
Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
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 DECL_POINTS_TO_READONLY
for VAR_DECLs.
* gimplify.cc (struct gimplify_omp_ctx):
Add 'hash_set<tree_operand_hash> *pt_readonly_ptrs' field.
(internal_get_tmp_var): Set
DECL_POINTS_TO_READONLY/SSA_NAME_POINTS_TO_READONLY_MEMORY for
new temp vars.
(build_omp_struct_comp_nodes):
Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
(gimplify_scan_omp_clauses): Collect OMP_CLAUSE_MAP_POINTS_TO_READONLY
to ctx->pt_readonly_ptrs.
* omp-low.cc (lower_omp_target): Set DECL_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.h (DECL_POINTS_TO_READONLY): New macro.
(OMP_CLAUSE_MAP_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.
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_))
