From 6e6cf7b08a65d22bd2b29ab97a3c01c4f646b21d Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Thu, 26 Jan 2017 12:05:03 +0100 Subject: [PATCH] re PR middle-end/79236 (Many libgomp tests fail if configured with --enable-offload-targets=nvptx-none but NVidia HW or libcuda.so.1 unavailable) 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 --- gcc/ChangeLog | 11 +++++++++++ gcc/omp-low.c | 25 ++++++++++++++++++++++--- 2 files changed, 33 insertions(+), 3 deletions(-) diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 8222331d35e..f485bb37255 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,14 @@ +2017-01-26 Jakub Jelinek + + 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 PR middle-end/79212 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 15209a3c354..ff0f4477cd7 100644 --- 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; } -- 2.30.2