openmp: Scale type precision of collapsed iterator variable

This sets the type precision of the collapsed iterator variable to the
sum of the precision of the collapsed loop variables, up to a maximum of
sizeof(long long) (i.e. 64-bits).

2021-03-01  Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/
	* omp-expand.cc (expand_oacc_for): Convert .tile variable to
	diff_type before multiplying.
	* omp-general.cc (omp_extract_for_data): Use accumulated precision
	of all collapsed for-loops as precision of iteration variable, up
	to the precision of a long long.

	libgomp/
	* testsuite/libgomp.c-c++-common/collapse-4.c: New.
	* testsuite/libgomp.fortran/collapse5.f90: New.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 1052b364689bdebf8b907e37e801a55c1f12fe24..712473c6378ec480984dec3a02750b049bc48474 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2021-03-01  Kwok Cheung Yeung  <kcy@codesourcery.com>
+
+       * omp-expand.cc (expand_oacc_for): Convert .tile variable to
+       diff_type before multiplying.
+       * omp-general.cc (omp_extract_for_data): Use accumulated precision
+       of all collapsed for-loops as precision of iteration variable, up
+       to the precision of a long long.
+
 2021-02-26  Andrew Stubbs  <ams@codesourcery.com>
 
        * dwarf2out.cc (gen_subprogram_die): Replace existing low/high PC

diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index e113231e52d4ca1911606ec70259b233b18b5040..316020772f2f47b8b367742cb15e59e932129c05 100644 (file)
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -7702,7 +7702,10 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd)
       tile_size = create_tmp_var (diff_type, ".tile_size");
       expr = build_int_cst (diff_type, 1);
       for (int ix = 0; ix < fd->collapse; ix++)
-       expr = fold_build2 (MULT_EXPR, diff_type, counts[ix].tile, expr);
+       {
+         tree tile = fold_convert (diff_type, counts[ix].tile);
+         expr = fold_build2 (MULT_EXPR, diff_type, tile, expr);
+       }
       expr = force_gimple_operand_gsi (&gsi, expr, true,
                                       NULL_TREE, true, GSI_SAME_STMT);
       ass = gimple_build_assign (tile_size, expr);

diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index a406c578f331c052d0f9c7cca6c4a5f8f068cfb5..8c7dce20dfb49c57ee503ff6055218e03f7c0822 100644 (file)
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -375,6 +375,7 @@ omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
          fd->non_rect = true;
        }
     }
+  int accum_iter_precision = 0;
   for (i = 0; i < cnt; i++)
     {
       if (i == 0
@@ -457,12 +458,28 @@ omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
        {
          if (fd->collapse == 1 && !fd->tiling)
            iter_type = TREE_TYPE (loop->v);
-         else if (i == 0
-                  || TYPE_PRECISION (iter_type)
-                     < TYPE_PRECISION (TREE_TYPE (loop->v)))
-           iter_type
-             = build_nonstandard_integer_type
-                 (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
+         else
+           {
+             int loop_precision = TYPE_PRECISION (TREE_TYPE (loop->v));
+             int iter_type_precision = 0;
+             const int max_accum_precision
+               = TYPE_PRECISION (long_long_unsigned_type_node);
+
+             accum_iter_precision += loop_precision;
+
+             if (i == 0
+                 || (loop_precision >= max_accum_precision
+                     && loop_precision >= TYPE_PRECISION (iter_type)))
+               iter_type_precision = loop_precision;
+             else if (TYPE_PRECISION (iter_type) < max_accum_precision)
+               iter_type_precision
+                 = MIN (1 << ceil_log2 (accum_iter_precision),
+                        max_accum_precision);
+
+             if (iter_type_precision)
+               iter_type = build_nonstandard_integer_type
+                             (iter_type_precision, 1);
+           }
        }
       else if (iter_type != long_long_unsigned_type_node)
        {

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 4bf75fd6861ac5ec66b95d4e378b7a2bee1a2b08..37795c1554e4d74e33652c75486ab902ecc69ae1 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-03-01  Kwok Cheung Yeung  <kcy@codesourcery.com>
+
+       * testsuite/libgomp.c-c++-common/collapse-4.c: New.
+       * testsuite/libgomp.fortran/collapse5.f90: New.
+
 2021-02-23  Andrew Stubbs  <ams@codesourcery.com>
 
        * plugin/plugin-nvptx.c (GOMP_OFFLOAD_alloc): Remove early call to

diff --git a/libgomp/testsuite/libgomp.c-c++-common/collapse-4.c b/libgomp/testsuite/libgomp.c-c++-common/collapse-4.c
new file mode 100644 (file)
index 0000000..c0af29f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/collapse-4.c
@@ -0,0 +1,23 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+  int i, j;
+  int count = 0;
+
+  #pragma omp parallel for collapse(2)
+    for (i = 0; i < 80000; i++)
+      for (j = 0; j < 80000; j++)
+       if (i == 66666 && j == 77777)
+         /* In the collapsed loop space, this is iteration
+            66666*80000+77777==5,333,357,777.  If the type of the iterator
+            for the collapsed loop is only a 32-bit unsigned int, then this
+            iteration will exceed its maximum range and be skipped.  */
+         count++;
+
+  if (count != 1)
+    abort ();
+}

diff --git a/libgomp/testsuite/libgomp.fortran/collapse5.f90 b/libgomp/testsuite/libgomp.fortran/collapse5.f90
new file mode 100644 (file)
index 0000000..5632d9b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/collapse5.f90
@@ -0,0 +1,23 @@
+! { dg-do run }
+
+program collapse5
+  implicit none
+
+  integer :: i, j
+  integer :: count = 0
+
+  !$omp parallel do collapse (2)
+    do i = 1, 80000
+      do j = 1, 80000
+        if (i .eq. 66666 .and. j .eq. 77777) then
+         ! In the collapsed loop space, this is iteration
+         ! 66666*80000+77777==5,333,357,777.  If the type of the iterator
+         ! for the collapsed loop is only a 32-bit unsigned int, then this
+         ! iteration will exceed its maximum range and be skipped.
+         count = count + 1
+       end if
+      end do
+    end do
+
+  if (count .ne. 1) stop 1
+end