re PR middle-end/79236 (Many libgomp tests fail if configured with --enable-offload...

	PR middle-end/79236
	* omp-low.c (struct omp_context): Add simt_stmt field.
	(scan_omp_for): Return omp_context *.
	(scan_omp_simd): Set simt_stmt on the non-_simt_ SIMD
	context to the _simt_ SIMD stmt.
	(lower_omp_for): For combined SIMD with sibling _simt_
	SIMD, make sure to use the same decls in _looptemp_
	clauses as in the sibling.

From-SVN: r244924

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 8222331d35eb4404d0d5dc6d258060e58ef36dd2..f485bb37255d482e10ca3f6c29932399cd782f2e 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,14 @@
+2017-01-26  Jakub Jelinek  <jakub@redhat.com>
+
+       PR middle-end/79236
+       * omp-low.c (struct omp_context): Add simt_stmt field.
+       (scan_omp_for): Return omp_context *.
+       (scan_omp_simd): Set simt_stmt on the non-_simt_ SIMD
+       context to the _simt_ SIMD stmt.
+       (lower_omp_for): For combined SIMD with sibling _simt_
+       SIMD, make sure to use the same decls in _looptemp_
+       clauses as in the sibling.
+
 2017-01-26  David Sherwood  <david.sherwood@arm.com>
 
        PR middle-end/79212

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 15209a3c3544efd3d30dd3db257a2f43a2299c7e..ff0f4477cd774f0e9c2eb6df14ffb16e942d770c 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -108,6 +108,10 @@ struct omp_context
      barriers should jump to during omplower pass.  */
   tree cancel_label;
 
+  /* The sibling GIMPLE_OMP_FOR simd with _simt_ clause or NULL
+     otherwise.  */
+  gimple *simt_stmt;
+
   /* What to do with variables with implicitly determined sharing
      attributes.  */
   enum omp_clause_default_kind default_kind;
@@ -2127,7 +2131,7 @@ check_oacc_kernel_gwv (gomp_for *stmt, omp_context *ctx)
 
 /* Scan a GIMPLE_OMP_FOR.  */
 
-static void
+static omp_context *
 scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
 {
   omp_context *ctx;
@@ -2200,6 +2204,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
       scan_omp_op (gimple_omp_for_incr_ptr (stmt, i), ctx);
     }
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
+  return ctx;
 }
 
 /* Duplicate #pragma omp simd, one for SIMT, another one for SIMD.  */
@@ -2241,7 +2246,7 @@ scan_omp_simd (gimple_stmt_iterator *gsi, gomp_for *stmt,
   gimple_bind_set_body (bind, seq);
   update_stmt (bind);
   scan_omp_for (new_stmt, outer_ctx);
-  scan_omp_for (stmt, outer_ctx);
+  scan_omp_for (stmt, outer_ctx)->simt_stmt = new_stmt;
 }
 
 /* Scan an OpenMP sections directive.  */
@@ -6750,11 +6755,15 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
        = (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR
           || gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP);
       tree outerc = NULL, *pc = gimple_omp_for_clauses_ptr (stmt);
+      tree simtc = NULL;
       tree clauses = *pc;
       if (taskreg_for)
        outerc
          = omp_find_clause (gimple_omp_taskreg_clauses (ctx->outer->stmt),
                             OMP_CLAUSE__LOOPTEMP_);
+      if (ctx->simt_stmt)
+       simtc = omp_find_clause (gimple_omp_for_clauses (ctx->simt_stmt),
+                                OMP_CLAUSE__LOOPTEMP_);
       for (i = 0; i < count; i++)
        {
          tree temp;
@@ -6767,12 +6776,22 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            }
          else
            {
-             temp = create_tmp_var (type);
+             /* If there are 2 adjacent SIMD stmts, one with _simt_
+                clause, another without, make sure they have the same
+                decls in _looptemp_ clauses, because the outer stmt
+                they are combined into will look up just one inner_stmt.  */
+             if (ctx->simt_stmt)
+               temp = OMP_CLAUSE_DECL (simtc);
+             else
+               temp = create_tmp_var (type);
              insert_decl_map (&ctx->outer->cb, temp, temp);
            }
          *pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_);
          OMP_CLAUSE_DECL (*pc) = temp;
          pc = &OMP_CLAUSE_CHAIN (*pc);
+         if (ctx->simt_stmt)
+           simtc = omp_find_clause (OMP_CLAUSE_CHAIN (simtc),
+                                    OMP_CLAUSE__LOOPTEMP_);
        }
       *pc = clauses;
     }