From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id B88CF386FC2B; Mon, 19 Apr 2021 10:37:17 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B88CF386FC2B MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r10-9716] OpenACC: Fix integer-type issue with collapse/tile [PR97880] X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/releases/gcc-10 X-Git-Oldrev: 67e95b6908b842e437089e867e5eed6dbbe264f2 X-Git-Newrev: af408874e3d94492ac08ba17462b3ff8ecb4d791 Message-Id: <20210419103717.B88CF386FC2B@sourceware.org> Date: Mon, 19 Apr 2021 10:37:17 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 19 Apr 2021 10:37:17 -0000 https://gcc.gnu.org/g:af408874e3d94492ac08ba17462b3ff8ecb4d791 commit r10-9716-gaf408874e3d94492ac08ba17462b3ff8ecb4d791 Author: Tobias Burnus Date: Fri Nov 27 11:17:50 2020 +0100 OpenACC: Fix integer-type issue with collapse/tile [PR97880] gcc/ChangeLog: PR c/97880 * omp-expand.c (expand_oacc_collapse_init, expand_oacc_collapse_vars): Use now passed diff_type. (expand_oacc_for): Take largest type for diff_type, taking tiling and collapsing into account. gcc/testsuite/ChangeLog: PR c/97880 * gcc.dg/goacc/tile-1.c: New test. (cherry picked from commit f324479caf0ac326534f4fcf72cb12991ccddb3d) Diff: --- gcc/omp-expand.c | 28 +++++++++++++++------------- gcc/testsuite/gcc.dg/goacc/tile-1.c | 10 ++++++++++ 2 files changed, 25 insertions(+), 13 deletions(-) diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index b73fcf0b0a2..af4059cc2ab 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -1548,8 +1548,8 @@ struct oacc_collapse static tree expand_oacc_collapse_init (const struct omp_for_data *fd, gimple_stmt_iterator *gsi, - oacc_collapse *counts, tree bound_type, - location_t loc) + oacc_collapse *counts, tree diff_type, + tree bound_type, location_t loc) { tree tiling = fd->tiling; tree total = build_int_cst (bound_type, 1); @@ -1566,17 +1566,12 @@ expand_oacc_collapse_init (const struct omp_for_data *fd, const omp_for_data_loop *loop = &fd->loops[ix]; tree iter_type = TREE_TYPE (loop->v); - tree diff_type = iter_type; tree plus_type = iter_type; gcc_assert (loop->cond_code == LT_EXPR || loop->cond_code == GT_EXPR); if (POINTER_TYPE_P (iter_type)) plus_type = sizetype; - if (POINTER_TYPE_P (diff_type) || TYPE_UNSIGNED (diff_type)) - diff_type = signed_type_for (diff_type); - if (TYPE_PRECISION (diff_type) < TYPE_PRECISION (integer_type_node)) - diff_type = integer_type_node; if (tiling) { @@ -1664,7 +1659,8 @@ expand_oacc_collapse_init (const struct omp_for_data *fd, static void expand_oacc_collapse_vars (const struct omp_for_data *fd, bool inner, gimple_stmt_iterator *gsi, - const oacc_collapse *counts, tree ivar) + const oacc_collapse *counts, tree ivar, + tree diff_type) { tree ivar_type = TREE_TYPE (ivar); @@ -1676,7 +1672,6 @@ expand_oacc_collapse_vars (const struct omp_for_data *fd, bool inner, const oacc_collapse *collapse = &counts[ix]; tree v = inner ? loop->v : collapse->outer; tree iter_type = TREE_TYPE (v); - tree diff_type = TREE_TYPE (collapse->step); tree plus_type = iter_type; enum tree_code plus_code = PLUS_EXPR; tree expr; @@ -1698,7 +1693,7 @@ expand_oacc_collapse_vars (const struct omp_for_data *fd, bool inner, } expr = fold_build2 (MULT_EXPR, diff_type, fold_convert (diff_type, expr), - collapse->step); + fold_convert (diff_type, collapse->step)); expr = fold_build2 (plus_code, iter_type, inner ? collapse->outer : collapse->base, fold_convert (plus_type, expr)); @@ -6039,6 +6034,12 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) plus_code = POINTER_PLUS_EXPR; plus_type = sizetype; } + for (int ix = fd->collapse; ix--;) + { + tree diff_type2 = TREE_TYPE (fd->loops[ix].step); + if (TYPE_PRECISION (diff_type) < TYPE_PRECISION (diff_type2)) + diff_type = diff_type2; + } if (POINTER_TYPE_P (diff_type) || TYPE_UNSIGNED (diff_type)) diff_type = signed_type_for (diff_type); if (TYPE_PRECISION (diff_type) < TYPE_PRECISION (integer_type_node)) @@ -6122,7 +6123,7 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) { gcc_assert (!gimple_in_ssa_p (cfun) && up); counts = XALLOCAVEC (struct oacc_collapse, fd->collapse); - tree total = expand_oacc_collapse_init (fd, &gsi, counts, + tree total = expand_oacc_collapse_init (fd, &gsi, counts, diff_type, TREE_TYPE (fd->loop.n2), loc); if (SSA_VAR_P (fd->loop.n2)) @@ -6284,7 +6285,7 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) gsi_insert_before (&gsi, ass, GSI_SAME_STMT); if (fd->collapse > 1 || fd->tiling) - expand_oacc_collapse_vars (fd, false, &gsi, counts, v); + expand_oacc_collapse_vars (fd, false, &gsi, counts, v, diff_type); if (fd->tiling) { @@ -6354,7 +6355,8 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) /* Initialize the user's loop vars. */ gsi = gsi_start_bb (elem_body_bb); - expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset); + expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset, + diff_type); } } diff --git a/gcc/testsuite/gcc.dg/goacc/tile-1.c b/gcc/testsuite/gcc.dg/goacc/tile-1.c new file mode 100644 index 00000000000..6898397ad5e --- /dev/null +++ b/gcc/testsuite/gcc.dg/goacc/tile-1.c @@ -0,0 +1,10 @@ +/* { dg-do compile } */ + +/* PR c/97880 */ + +void f () +{ + #pragma acc parallel loop tile(2, 3) + for (int i = 0; i < 8; i++) + for (long j = 0; j < 8; j++); +}