From patchwork Wed Dec 15 15:54:19 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48954 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id EC6E6385E82B for ; Wed, 15 Dec 2021 16:03:05 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 177363858039 for ; Wed, 15 Dec 2021 15:55:49 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 177363858039 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: FkXb6U0IGdeGzaR1J0RGJrb5PC8qaZNaS4mLK/Bp73rLbRFHhDZvH/m+LMqOsfmyKEtCKevgOu IlTOBFStci9/kjNFrpyjxCdw6sLheXYQnnh1XRlKnai1yYm+eOD2GLfb29IGnCBBJxBsolTihk Vv+7m56eO+GBN9NMhtEHGjwdKT34GZZ3BWDwnkgqoY9KotU5wpeJLIRisVL3icesU4OBEyKt3K agAKF4cSuFFbv0lQeTalKv4QRyftev4jDX9z0g5NTrAssqK8qM+pleZxBh7TwddDzhYX5I4XoK AOJwVt5Jo7cqD/h8Dav46/ix X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69736557" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:48 -0800 IronPort-SDR: WzDbXTaVXkzrTKS9q/BTZ9etubChrTz1+fij98IzNJ6UaWU9ZixyaR9lNtUUhzkuy/SUV2J5AH /JPeSQZZinVU7rI2M1MscyTiDoB85igZB6uXK4oey6LFS0TCal4oSrEQZeWoyufAWpMhkoLvCW epeEX1k06nbF5SchHpXUmAGrSYmKeJhOb81FFRE7CvssCokRV63f1Ig/A6wyjJ3DuokQeNYC3G yJ/1vJjqIaiNBNieZNwdAwJHL8FXu79VF8EzlpGsHM0HeyRO7nvuWwDa8SaZhQpJjeWcgm71ds kdU= From: Frederik Harwath To: Subject: [PATCH 12/40] Relax some restrictions on the loop bound in kernels loop annotation. Date: Wed, 15 Dec 2021 16:54:19 +0100 Message-ID: <20211215155447.19379-13-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Sandra Loosemore , thomas@codesourcery.com, joseph@codesourcery.com, nathan@acm.org Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore 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 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/c-omp.c | 120 ++++++++++++++++-- .../goacc/kernels-loop-annotation-21.c | 42 ++++++ .../goacc/kernels-loop-annotation-22.c | 41 ++++++ 3 files changed, 194 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 -- 2.33.0 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index e73fb5d01f7e..dc63d304ca67 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -3165,6 +3165,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 @@ -3333,16 +3443,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/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 000000000000..f87444ede4b4 --- /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 000000000000..6a5099d2ff9d --- /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" } } */