&& !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
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
--- /dev/null
+/* { 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" } } */
--- /dev/null
+/* { 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" } } */
projects / thirdparty / gcc.git / commitdiff
