[gcc r10-9716] OpenACC: Fix integer-type issue with collapse/tile [PR97880]

[gcc r10-9716] OpenACC: Fix integer-type issue with collapse/tile [PR97880]

[Tobias Burnus]

https://gcc.gnu.org/g:af408874e3d94492ac08ba17462b3ff8ecb4d791

commit r10-9716-gaf408874e3d94492ac08ba17462b3ff8ecb4d791
Author: Tobias Burnus <tobias@codesourcery.com>
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++);
+}