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

2 files changed, 77 insertions, 0 deletions

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  <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 00000000000..84ad55823cb
--- /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;
+}