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

From-SVN: r218158

4 files changed, 97 insertions, 10 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9696664..3165e49 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -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
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 3924282..6c5774c 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 15e5ae3..7ec3fb1 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -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
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 0000000..84ad558
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-critical-1.c
@@ -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;
+}