public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc(refs/vendors/redhat/heads/gcc-8-branch)] modulo-sched: skip loops with strange register defs [PR100225]
@ 2021-05-14 14:57 Jakub Jelinek
0 siblings, 0 replies; only message in thread
From: Jakub Jelinek @ 2021-05-14 14:57 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:d43be80f8d0f2bd9f2062c51cee0d0766a7f22d3
commit d43be80f8d0f2bd9f2062c51cee0d0766a7f22d3
Author: Roman Zhuykov <zhroma@ispras.ru>
Date: Thu May 6 15:01:37 2021 +0300
modulo-sched: skip loops with strange register defs [PR100225]
PR84878 fix adds an assertion which can fail, e.g. when stack pointer
is adjusted inside the loop. We have to prevent it and search earlier
for any 'strange' instruction. The solution is to skip the whole loop
if using 'note_stores' we found that one of hard registers is in
'df->regular_block_artificial_uses' set.
Also patch properly prohibit not single-set instruction in loop body.
gcc/ChangeLog:
PR rtl-optimization/100225
PR rtl-optimization/84878
* modulo-sched.c (sms_schedule): Use note_stores to skip loops
where we have an instruction which touches (writes) any hard
register from df->regular_block_artificial_uses set.
Allow not-single-set instruction only right before basic block
tail.
gcc/testsuite/ChangeLog:
PR rtl-optimization/100225
PR rtl-optimization/84878
* gcc.dg/pr100225.c: New test.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c: New test.
(cherry picked from commit 4cf3b10f27b1994cf4a9eb12079d85412ebc7cad)
Diff:
---
gcc/modulo-sched.c | 56 +-
gcc/testsuite/gcc.dg/pr100225.c | 15 +
.../libgomp.oacc-c-c++-common/atomic_capture-3.c | 1627 ++++++++++++++++++++
3 files changed, 1677 insertions(+), 21 deletions(-)
diff --git a/gcc/modulo-sched.c b/gcc/modulo-sched.c
index d8e6c9309e1..40395160a5b 100644
--- a/gcc/modulo-sched.c
+++ b/gcc/modulo-sched.c
@@ -44,6 +44,7 @@ along with GCC; see the file COPYING3. If not see
#include "tree-pass.h"
#include "dbgcnt.h"
#include "loop-unroll.h"
+#include "hard-reg-set.h"
#ifdef INSN_SCHEDULING
@@ -1358,6 +1359,7 @@ sms_schedule (void)
basic_block condition_bb = NULL;
edge latch_edge;
HOST_WIDE_INT trip_count, max_trip_count;
+ HARD_REG_SET prohibited_regs;
loop_optimizer_init (LOOPS_HAVE_PREHEADERS
| LOOPS_HAVE_RECORDED_EXITS);
@@ -1387,6 +1389,8 @@ sms_schedule (void)
We use loop->num as index into this array. */
g_arr = XCNEWVEC (ddg_ptr, number_of_loops (cfun));
+ REG_SET_TO_HARD_REG_SET (prohibited_regs, &df->regular_block_artificial_uses);
+
if (dump_file)
{
fprintf (dump_file, "\n\nSMS analysis phase\n");
@@ -1475,23 +1479,31 @@ sms_schedule (void)
}
/* Don't handle BBs with calls or barriers
- or !single_set with the exception of instructions that include
- count_reg---these instructions are part of the control part
- that do-loop recognizes.
+ or !single_set with the exception of do-loop control part insns.
??? Should handle insns defining subregs. */
- for (insn = head; insn != NEXT_INSN (tail); insn = NEXT_INSN (insn))
- {
- rtx set;
-
- if (CALL_P (insn)
- || BARRIER_P (insn)
- || (NONDEBUG_INSN_P (insn) && !JUMP_P (insn)
- && !single_set (insn) && GET_CODE (PATTERN (insn)) != USE
- && !reg_mentioned_p (count_reg, insn))
- || (INSN_P (insn) && (set = single_set (insn))
- && GET_CODE (SET_DEST (set)) == SUBREG))
- break;
- }
+ for (insn = head; insn != NEXT_INSN (tail); insn = NEXT_INSN (insn))
+ {
+ if (INSN_P (insn))
+ {
+ HARD_REG_SET regs;
+ CLEAR_HARD_REG_SET (regs);
+ note_stores (PATTERN (insn), record_hard_reg_sets, ®s);
+ if (hard_reg_set_intersect_p (regs, prohibited_regs))
+ break;
+ }
+
+ if (CALL_P (insn)
+ || BARRIER_P (insn)
+ || (INSN_P (insn) && single_set (insn)
+ && GET_CODE (SET_DEST (single_set (insn))) == SUBREG)
+ /* Not a single set. */
+ || (NONDEBUG_INSN_P (insn) && !JUMP_P (insn)
+ && !single_set (insn) && GET_CODE (PATTERN (insn)) != USE
+ /* But non-single-set allowed in one special case. */
+ && (insn != prev_nondebug_insn (tail)
+ || !reg_mentioned_p (count_reg, insn))))
+ break;
+ }
if (insn != NEXT_INSN (tail))
{
@@ -1501,11 +1513,13 @@ sms_schedule (void)
fprintf (dump_file, "SMS loop-with-call\n");
else if (BARRIER_P (insn))
fprintf (dump_file, "SMS loop-with-barrier\n");
- else if ((NONDEBUG_INSN_P (insn) && !JUMP_P (insn)
- && !single_set (insn) && GET_CODE (PATTERN (insn)) != USE))
- fprintf (dump_file, "SMS loop-with-not-single-set\n");
- else
- fprintf (dump_file, "SMS loop with subreg in lhs\n");
+ else if (INSN_P (insn) && single_set (insn)
+ && GET_CODE (SET_DEST (single_set (insn))) == SUBREG)
+ fprintf (dump_file, "SMS loop with subreg in lhs\n");
+ else
+ fprintf (dump_file,
+ "SMS loop-with-not-single-set-or-prohibited-reg\n");
+
print_rtl_single (dump_file, insn);
}
diff --git a/gcc/testsuite/gcc.dg/pr100225.c b/gcc/testsuite/gcc.dg/pr100225.c
new file mode 100644
index 00000000000..b32163441a3
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/pr100225.c
@@ -0,0 +1,15 @@
+/* PR rtl-optimization/100225 */
+/* { dg-do compile } */
+/* { dg-options "-O1 -fmodulo-sched" } */
+
+void
+vorbis_synthesis_lapout (void);
+
+void
+ov_info (int **lappcm, int ov_info_i)
+{
+ while (ov_info_i < 1)
+ lappcm[ov_info_i++] = __builtin_alloca (1);
+
+ vorbis_synthesis_lapout ();
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
new file mode 100644
index 00000000000..b976094998f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
@@ -0,0 +1,1627 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fmodulo-sched -fmodulo-sched-allow-regmoves" } */
+
+#include <stdlib.h>
+
+int
+main(int argc, char **argv)
+{
+ int iexp, igot, imax, imin;
+ long long lexp, lgot;
+ int N = 32;
+ int i;
+ int idata[N];
+ long long ldata[N];
+ float fexp, fgot;
+ float fdata[N];
+
+ igot = 1234;
+ iexp = 31;
+
+ for (i = 0; i < N; i++)
+ idata[i] = i;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { idata[i] = igot; igot = i; }
+ }
+
+ imax = 0;
+ imin = N;
+
+ for (i = 0; i < N; i++)
+ {
+ imax = idata[i] > imax ? idata[i] : imax;
+ imin = idata[i] < imin ? idata[i] : imin;
+ }
+
+ if (imax != 1234 || imin != 0)
+ abort ();
+
+ return 0;
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { idata[i] = igot; igot++; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { idata[i] = igot; ++igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { ++igot; idata[i] = igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { igot++; idata[i] = igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { idata[i] = igot; igot--; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { idata[i] = igot; --igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { --igot; idata[i] = igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+#pragma acc atomic capture
+ { igot--; idata[i] = igot; }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ /* BINOP = + */
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot += expr; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot += expr; idata[i] = igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot = igot + expr; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot = expr + igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot = igot + expr; idata[i] = igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+
+ igot = 0;
+ iexp = 32;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot = expr + igot; idata[i] = igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ /* BINOP = * */
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot *= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { lgot *= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot * expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr * lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { lgot = lgot * expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2;
+
+#pragma acc atomic capture
+ { lgot = expr * lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = - */
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot -= expr; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot -= expr; idata[i] = igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 32;
+ iexp = 0;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot = igot - expr; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 1;
+ iexp = 1;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { idata[i] = igot; igot = expr - igot; }
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ if (i % 2 == 0)
+ {
+ if (idata[i] != 1)
+ abort ();
+ }
+ else
+ {
+ if (idata[i] != 0)
+ abort ();
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 1;
+ iexp = -31;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot = igot - expr; idata[i] = igot; }
+ }
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ igot = 1;
+ iexp = 1;
+
+#pragma acc data copy (igot, idata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ int expr = 1;
+
+#pragma acc atomic capture
+ { igot = expr - igot; idata[i] = igot; }
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ if (i % 2 == 0)
+ {
+ if (idata[i] != 0)
+ abort ();
+ }
+ else
+ {
+ if (idata[i] != 1)
+ abort ();
+ }
+
+ if (iexp != igot)
+ abort ();
+
+ /* BINOP = / */
+ lgot = 1LL << 32;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot /= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL << 32;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { lgot /= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL << 32;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot / expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 2LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr / lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 2LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { lgot = lgot / expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 2LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { lgot = expr / lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = & */
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot &= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ iexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot &= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot & expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr & lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ iexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = lgot & expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = expr & lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = ^ */
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1 << i;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot ^= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ iexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot ^= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot ^ expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr ^ lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ iexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = lgot ^ expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = ~0LL;
+ lexp = 0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = expr ^ lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = | */
+ lgot = 0LL;
+ lexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1 << i;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot |= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 0LL;
+ iexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot |= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 0LL;
+ lexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot | expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 0LL;
+ lexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr | lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 0LL;
+ iexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = lgot | expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 0LL;
+ lexp = ~0LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = ~(1 << i);
+
+#pragma acc atomic capture
+ { lgot = expr | lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = << */
+ lgot = 1LL;
+ lexp = 1LL << N;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot <<= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ iexp = 1LL << N;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { lgot <<= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << N;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot << expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < 1; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr << lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < 1; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { lgot = lgot << expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 2LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < 1; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { lgot = expr << lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ /* BINOP = >> */
+ lgot = 1LL << N;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot >>= expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL << N;
+ iexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { lgot >>= expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL << N;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = lgot >> expr; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << (N - 1);
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < 1; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { ldata[i] = lgot; lgot = expr >> lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL << N;
+ lexp = 1LL;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL;
+
+#pragma acc atomic capture
+ { lgot = lgot >> expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << (N - 1);
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < 1; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { lgot = expr >> lgot; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ // FLOAT FLOAT FLOAT
+
+ /* BINOP = + */
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot += expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot += expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { idata[i] = fgot; fgot = fgot + expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = expr + fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot = fgot + expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 0.0;
+ fexp = 32.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot = expr + fgot; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ /* BINOP = * */
+ fgot = 1.0;
+ fexp = 8192.0*8192.0*64.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot *= expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 8192.0*8192.0*64.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fgot *= expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 8192.0*8192.0*64.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = fgot * expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 8192.0*8192.0*64.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = expr * fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ lgot = 1LL;
+ lexp = 1LL << 32;
+
+#pragma acc data copy (lgot, ldata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2LL;
+
+#pragma acc atomic capture
+ { lgot = lgot * expr; ldata[i] = lgot; }
+ }
+ }
+
+ if (lexp != lgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 8192.0*8192.0*64.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 2;
+
+#pragma acc atomic capture
+ { fgot = expr * fgot; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ /* BINOP = - */
+ fgot = 32.0;
+ fexp = 0.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot -= expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 32.0;
+ fexp = 0.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot -= expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 32.0;
+ fexp = 0.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = fgot - expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = expr - fgot; }
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ if (i % 2 == 0)
+ {
+ if (fdata[i] != 1.0)
+ abort ();
+ }
+ else
+ {
+ if (fdata[i] != 0.0)
+ abort ();
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = -31.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot = fgot - expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 1.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fgot = expr - fgot; fdata[i] = fgot; }
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ if (i % 2 == 0)
+ {
+ if (fdata[i] != 0.0)
+ abort ();
+ }
+ else
+ {
+ if (fdata[i] != 1.0)
+ abort ();
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ /* BINOP = / */
+ fgot = 8192.0*8192.0*64.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot /= expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 8192.0*8192.0*64.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fgot /= expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 8192.0*8192.0*64.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = fgot / expr; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 8192.0*8192.0*64.0;
+ fexp = 1.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 1.0;
+
+#pragma acc atomic capture
+ { fdata[i] = fgot; fgot = expr / fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 4.0;
+ fexp = 4.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ long long expr = 1LL << N;
+
+#pragma acc atomic capture
+ { fgot = fgot / expr; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ fgot = 4.0;
+ fexp = 4.0;
+
+#pragma acc data copy (fgot, fdata[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ {
+ float expr = 2.0;
+
+#pragma acc atomic capture
+ { fgot = expr / fgot; fdata[i] = fgot; }
+ }
+ }
+
+ if (fexp != fgot)
+ abort ();
+
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-14 14:57 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-14 14:57 [gcc(refs/vendors/redhat/heads/gcc-8-branch)] modulo-sched: skip loops with strange register defs [PR100225] Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).