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

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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9696664645f59bebba06eaa2d76798a088a27ed8..3165e49961984350ba641f65037502032a457030 100644 (file)
--- 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 3924282cae248b4490f1f5f284bbf5c004c6dc5e..6c5774c2606f39d7aa7b26333171dac2e07ac0e3 100644 (file)
--- 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 15e5ae35b41f542d24922b6cc303b62eda4949fd..7ec3fb1e261956b149d9901059bab4f0a599f9ea 100644 (file)
--- 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 (file)
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;
+}