[OG10] [committed 2/2] Relax some restrictions on the loop bound in kernels loop annotation.
Sandra Loosemore
sandra@codesourcery.com
Sun Aug 30 19:46:11 GMT 2020
OpenACC loop semantics require that the loop bound be computable
before entering the loop, rather than the C/C++ semantics where the
end test is evaluated on every iteration. Formerly the kernels loop
annotater permitted only constants and variables not modified in the
loop body in the loop bound expression. This patch relaxes those
restrictions somewhat to allow many forms of expressions involving
such constants and variables, including calls to constant functions.
2020-08-30 Sandra Loosemore <sandra@codesourcery.com>
gcc/c-family/
* c-omp.c (end_test_ok_for_annotation_r): New.
(end_test_ok_for_annotation): New.
(check_and_annotate_for_loop): Use the new helper function.
gcc/testsuite/
* c-c++-common/goacc/kernels-loop-annotation-21.c: New.
* c-c++-common/goacc/kernels-loop-annotation-22.c: New.
---
gcc/c-family/ChangeLog.omp | 6 ++
gcc/c-family/c-omp.c | 120 +++++++++++++++++++--
gcc/testsuite/ChangeLog.omp | 5 +
.../goacc/kernels-loop-annotation-21.c | 42 ++++++++
.../goacc/kernels-loop-annotation-22.c | 41 +++++++
5 files changed, 205 insertions(+), 9 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c
diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index a57bcef..ecec040 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,4 +1,10 @@
2020-08-30 Sandra Loosemore <sandra@codesourcery.com>
+
+ * c-omp.c (end_test_ok_for_annotation_r): New.
+ (end_test_ok_for_annotation): New.
+ (check_and_annotate_for_loop): Use the new helper function.
+
+2020-08-30 Sandra Loosemore <sandra@codesourcery.com>
* c-omp.c (annotate_for_loop): Move initializer processing...
(check_and_annotate_for_loop): ... to here. Allow the loop
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 79f4bef..b9531cc 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2470,6 +2470,116 @@ is_local_var (tree decl)
&& !TREE_ADDRESSABLE (decl));
}
+/* EXP is a loop bound expression for a comparison against local
+ variable DECL. Check whether this is potentially valid in an OpenACC loop
+ context, namely that it can be precomputed when entering the loop
+ construct per the OpenACC specification. Local variables referenced
+ in both DECL and EXP that may not be modified in the body of the loop
+ are added to the list in INFO to be checked later.
+
+ FIXME: Ideally we would like to make this test permissive rather than
+ restrictive, and allow the later conversion of the "auto" attribute to
+ either "seq" or "independent" to make the determination using dataflow,
+ alias analysis, etc rather than a tree traversal. But presently it does
+ not do that and always just hoists the loop bound expression. So the
+ current implementation only considers expressions involving unmodified
+ local variables and constants, using a tree walk. */
+
+static tree
+end_test_ok_for_annotation_r (tree *tp, int *walk_subtrees,
+ void *data)
+{
+ tree exp = *tp;
+ struct annotation_info *info = (struct annotation_info *) data;
+
+ switch (TREE_CODE_CLASS (TREE_CODE (exp)))
+ {
+ case tcc_constant:
+ /* Constants are trivially known to be invariant. */
+ return NULL_TREE;
+
+ case tcc_declaration:
+ if (is_local_var (exp))
+ {
+ tree t;
+ /* Add it to the list of variables that can't be modified in the
+ loop, only if not already present. */
+ for (t = info->vars; t && TREE_VALUE (t) != exp;
+ t = TREE_CHAIN (t))
+ ;
+ if (!t)
+ info->vars = tree_cons (NULL_TREE, exp, info->vars);
+ return NULL_TREE;
+ }
+ else if (TREE_CODE (exp) == VAR_DECL && TREE_READONLY (exp))
+ return NULL_TREE;
+ else if (TREE_CODE (exp) == FUNCTION_DECL)
+ return NULL_TREE;
+ break;
+
+ case tcc_unary:
+ case tcc_binary:
+ case tcc_comparison:
+ /* Allow arithmetic expressions and comparisons provided
+ that the operands are good. */
+ return NULL_TREE;
+
+ default:
+ /* Handle some special cases. */
+ switch (TREE_CODE (exp))
+ {
+ case COND_EXPR:
+ case TRUTH_ANDIF_EXPR:
+ case TRUTH_ORIF_EXPR:
+ case TRUTH_AND_EXPR:
+ case TRUTH_OR_EXPR:
+ case TRUTH_XOR_EXPR:
+ case TRUTH_NOT_EXPR:
+ /* ?: and boolean operators are OK. */
+ return NULL_TREE;
+
+ case CALL_EXPR:
+ /* Allow calls to constant functions with invariant operands. */
+ {
+ tree fndecl = get_callee_fndecl (exp);
+ if (fndecl && TREE_READONLY (fndecl))
+ return NULL_TREE;
+ }
+ break;
+
+ case ADDR_EXPR:
+ /* We can expect addresses of things to be invariant. */
+ return NULL_TREE;
+
+ default:
+ break;
+ }
+ }
+
+ /* Reject anything else. */
+ *walk_subtrees = 0;
+ return exp;
+}
+
+static bool
+end_test_ok_for_annotation (tree decl, tree exp,
+ struct annotation_info *info)
+{
+ /* Traversal returns NULL_TREE if all is well. */
+ if (!walk_tree (&exp, end_test_ok_for_annotation_r, info, NULL))
+ {
+ /* So far, so good. Check the decl against any variables collected
+ in the exp. */
+ tree t;
+ for (t = info->vars; t; t = TREE_CHAIN (t))
+ if (TREE_VALUE (t) == decl)
+ return false;
+ info->vars = tree_cons (NULL_TREE, decl, info->vars);
+ return true;
+ }
+ return false;
+}
+
/* The initializer for a FOR_STMT is sometimes wrapped in various other
language-specific tree structures. We need a hook to unwrap them.
This function takes a tree argument and should return either a
@@ -2638,16 +2748,8 @@ check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi,
limit_exp = TREE_OPERAND (cond, 0);
if (!limit_exp
- || (!is_local_var (limit_exp)
- && (TREE_CODE_CLASS (TREE_CODE (limit_exp)) != tcc_constant)))
+ || !end_test_ok_for_annotation (decl, limit_exp, &loop_info))
do_not_annotate_loop (&loop_info, as_invalid_predicate, cond);
- else
- {
- /* These variables must not be assigned to in the loop. */
- loop_info.vars = tree_cons (NULL_TREE, decl, loop_info.vars);
- if (TREE_CODE_CLASS (TREE_CODE (limit_exp)) != tcc_constant)
- loop_info.vars = tree_cons (NULL_TREE, limit_exp, loop_info.vars);
- }
}
/* Walk the body. This will process any nested loops, so we have to do it
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index c0d7e48..dd54db2 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2020-08-30 Sandra Loosemore <sandra@codesourcery.com>
+
+ * c-c++-common/goacc/kernels-loop-annotation-21.c: New.
+ * c-c++-common/goacc/kernels-loop-annotation-22.c: New.
+
2020-08-25 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c
new file mode 100644
index 0000000..f87444e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c
@@ -0,0 +1,42 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test for rejecting annotation on loops that have various subexpressions
+ in the loop end test that are not loop-invariant. */
+
+extern int g (int);
+extern int x;
+extern int gg (int, int) __attribute__ ((const));
+
+void f (float *a, float *b, int n)
+{
+
+ int j;
+#pragma acc kernels
+ {
+ /* Non-constant function call. */
+ for (int i = 0; i < g(n); i++) /* { dg-warning "loop cannot be annotated" } */
+ a[i] = b[i];
+
+ /* Global variable. */
+ for (int i = x; i < n + x; i++) /* { dg-warning "loop cannot be annotated" } */
+ a[i] = b[i];
+
+ /* Explicit reference to the loop variable. */
+ for (int i = 0; i < gg (i, n); i++) /* { dg-warning "loop cannot be annotated" } */
+ a[i] = b[i];
+
+ /* Reference to a variable that is modified in the body of the loop. */
+ j = 0;
+ for (int i = 0; i < gg (j, n); i++) /* { dg-warning "loop cannot be annotated" } */
+ {
+ a[i] = b[i];
+ j = i;
+ }
+
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c
new file mode 100644
index 0000000..6a5099d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c
@@ -0,0 +1,41 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test for accepting annotation on loops that have various forms of
+ loop-invariant expressions in their end test. */
+
+extern const int x;
+extern int g (int) __attribute__ ((const));
+
+void f (float *a, float *b, int n)
+{
+
+ int j;
+#pragma acc kernels
+ {
+ /* Reversed form of comparison. */
+ for (int i = 0; n >= i; i++)
+ a[i] = b[i];
+
+ /* Constant function call. */
+ for (int i = 0; i < g(n); i++)
+ a[i] = b[i];
+
+ /* Constant global variable. */
+ for (int i = 0; i < x; i++)
+ a[i] = b[i];
+
+ /* Complicated expression involving conditionals, etc. */
+ for (int i = 0; i < ((x == 4) ? (n << 2) : (n << 3)); i++)
+ a[i] = b[i];
+
+ /* Reference to a local variable not modified in the loop. */
+ j = ((x == 4) ? (n << 2) : (n << 3));
+ for (int i = 0; i < j; i++)
+ a[i] = b[i];
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 5 "original" } } */
--
2.8.1
More information about the Gcc-patches
mailing list