https://gcc.gnu.org/g:22da577a9f974b499e4364719086e4835dfbd529
commit 22da577a9f974b499e4364719086e4835dfbd529
Author: Kwok Cheung Yeung <k...@codesourcery.com>
Date: Mon Mar 1 14:15:30 2021 -0800
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 <k...@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:
---
gcc/ChangeLog.omp | 8 +++++
gcc/omp-expand.cc | 5 ++-
gcc/omp-general.cc | 39 +++++++++++++++-------
libgomp/ChangeLog.omp | 5 +++
.../testsuite/libgomp.c-c++-common/collapse-4.c | 23 +++++++++++++
libgomp/testsuite/libgomp.fortran/collapse5.f90 | 23 +++++++++++++
6 files changed, 90 insertions(+), 13 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d000c3614a8..8199cacbdd7 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2021-03-01 Kwok Cheung Yeung <k...@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-01 Chung-Lin Tang <clt...@codesourcery.com>
* omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE'
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index a8782a09df8..98b9b44e9d0 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -7778,7 +7778,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 2c095200d5b..9a125a28afa 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -395,6 +395,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
@@ -478,19 +479,33 @@ 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)))
+ else
{
- if (TREE_CODE (iter_type) == BITINT_TYPE
- || TREE_CODE (TREE_TYPE (loop->v)) == BITINT_TYPE)
- iter_type
- = build_bitint_type (TYPE_PRECISION (TREE_TYPE (loop->v)),
- 1);
- else
- iter_type
- = build_nonstandard_integer_type
- (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
+ 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)
+ {
+ if (TREE_CODE (iter_type) == BITINT_TYPE
+ || TREE_CODE (TREE_TYPE (loop->v)) == BITINT_TYPE)
+ iter_type = build_bitint_type (iter_type_precision, 1);
+ else
+ 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 5bf3a8a6890..6931529c366 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-03-01 Kwok Cheung Yeung <k...@codesourcery.com>
+
+ * testsuite/libgomp.c-c++-common/collapse-4.c: New.
+ * testsuite/libgomp.fortran/collapse5.f90: New.
+
2021-08-03 Andrew Stubbs <a...@codesourcery.com>
* config/gcn/bar.h (gomp_barrier_init): Limit thread count to the
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
index 00000000000..c0af29f5463
--- /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
index 00000000000..5632d9bab02
--- /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
