mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-14 09:10:48 +08:00
Code Issues Packages Projects Releases Wiki Activity

omp-low.c (lower_omp_critical): Mark critical sections inside target functions as offloadable.

Browse Source 
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 <ilya.verbin@intel.com>

From-SVN: r218158
This commit is contained in:
Andrey Turetskiy 2014-11-28 13:59:49 +00:00 committed by Ilya Verbin
parent f672337f6d
commit 476ff78736
4 changed files with 97 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
gcc/ChangeLog
View File

@ -1,3 +1,9 @@
2014-11-28 Andrey Turetskiy <andrey.turetskiy@intel.com>
Ilya Verbin <ilya.verbin@intel.com>
* omp-low.c (lower_omp_critical): Mark critical sections
inside target functions as offloadable.
2014-11-28 Ilya Verbin <ilya.verbin@intel.com>
* lto-wrapper.c (run_gcc): Set have_lto and have_offload if at least one

24
gcc/omp-low.c
View File

@ -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));

5
libgomp/ChangeLog
View File

@ -1,3 +1,8 @@
2014-11-28 Andrey Turetskiy <andrey.turetskiy@intel.com>
Ilya Verbin <ilya.verbin@intel.com>
* testsuite/libgomp.c/target-critical-1.c: New test.
2014-11-26 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/examples-4/e.53.4.c: Add -DITESTITERS=20

72
libgomp/testsuite/libgomp.c/target-critical-1.c Normal file
View File

@ -0,0 +1,72 @@
/* { dg-do run } */
#include <omp.h>
#include <stdlib.h>
#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;
}