From 476ff787366cedb6f4123c0a5647522698570d5f Mon Sep 17 00:00:00 2001 From: Andrey Turetskiy Date: Fri, 28 Nov 2014 13:59:49 +0000 Subject: [PATCH] omp-low.c (lower_omp_critical): Mark critical sections inside target functions as offloadable. gcc/ * omp-low.c (lower_omp_critical): Mark critical sections inside target functions as offloadable. libgomp/ * testsuite/libgomp.c/target-critical-1.c: New test. Co-Authored-By: Ilya Verbin From-SVN: r218158 --- gcc/ChangeLog | 6 ++ gcc/omp-low.c | 24 ++++--- libgomp/ChangeLog | 5 ++ .../testsuite/libgomp.c/target-critical-1.c | 72 +++++++++++++++++++ 4 files changed, 97 insertions(+), 10 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/target-critical-1.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9696664645f..3165e499619 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2014-11-28 Andrey Turetskiy + Ilya Verbin + + * omp-low.c (lower_omp_critical): Mark critical sections + inside target functions as offloadable. + 2014-11-28 Ilya Verbin * lto-wrapper.c (run_gcc): Set have_lto and have_offload if at least one diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 3924282cae2..6c5774c2606 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -9366,16 +9366,6 @@ lower_omp_critical (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_ARTIFICIAL (decl) = 1; DECL_IGNORED_P (decl) = 1; - /* If '#pragma omp critical' is inside target region, the symbol must - be marked for offloading. */ - omp_context *octx; - for (octx = ctx->outer; octx; octx = octx->outer) - if (is_targetreg_ctx (octx)) - { - varpool_node::get_create (decl)->offloadable = 1; - break; - } - varpool_node::finalize_decl (decl); critical_name_mutexes->put (name, decl); @@ -9383,6 +9373,20 @@ lower_omp_critical (gimple_stmt_iterator *gsi_p, omp_context *ctx) else decl = *n; + /* If '#pragma omp critical' is inside target region or + inside function marked as offloadable, the symbol must be + marked as offloadable too. */ + omp_context *octx; + if (cgraph_node::get (current_function_decl)->offloadable) + varpool_node::get_create (decl)->offloadable = 1; + else + for (octx = ctx->outer; octx; octx = octx->outer) + if (is_targetreg_ctx (octx)) + { + varpool_node::get_create (decl)->offloadable = 1; + break; + } + lock = builtin_decl_explicit (BUILT_IN_GOMP_CRITICAL_NAME_START); lock = build_call_expr_loc (loc, lock, 1, build_fold_addr_expr_loc (loc, decl)); diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 15e5ae35b41..7ec3fb1e261 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2014-11-28 Andrey Turetskiy + Ilya Verbin + + * testsuite/libgomp.c/target-critical-1.c: New test. + 2014-11-26 Jakub Jelinek * testsuite/libgomp.c/examples-4/e.53.4.c: Add -DITESTITERS=20 diff --git a/libgomp/testsuite/libgomp.c/target-critical-1.c b/libgomp/testsuite/libgomp.c/target-critical-1.c new file mode 100644 index 00000000000..84ad55823cb --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-critical-1.c @@ -0,0 +1,72 @@ +/* { dg-do run } */ + +#include +#include + +#define N 2000 + +#pragma omp declare target +int foo () +{ + int A[N]; + int i, nthreads; + int res = 0; + + #pragma omp parallel shared (A, nthreads) + { + #pragma omp master + nthreads = omp_get_num_threads (); + + #pragma omp for + for (i = 0; i < N; i++) + A[i] = 0; + + #pragma omp critical (crit1) + for (i = 0; i < N; i++) + A[i]++; + } + + for (i = 0; i < N; i++) + if (A[i] != nthreads) + res = 1; + + return res; +} +#pragma omp end declare target + +int main () +{ + int res1, res2; + + #pragma omp target map (from: res1, res2) + { + int B[N]; + int i, nthreads; + + res1 = foo (); + + #pragma omp parallel shared (B, nthreads) + { + #pragma omp master + nthreads = omp_get_num_threads (); + + #pragma omp for + for (i = 0; i < N; i++) + B[i] = 0; + + #pragma omp critical (crit2) + for (i = 0; i < N; i++) + B[i]++; + } + + res2 = 0; + for (i = 0; i < N; i++) + if (B[i] != nthreads) + res2 = 1; + } + + if (res1 || res2) + abort (); + + return 0; +} -- 2.30.2