diff mbox

[OpenMP] Fix named critical sections inside target functions

Message ID 20141127235211.GC30075@msticlxl57.ims.intel.com
State New
Headers show

Commit Message

Ilya Verbin Nov. 27, 2014, 11:52 p.m. UTC 
On 21 Nov 21:36, Jakub Jelinek wrote:
> On Fri, Nov 21, 2014 at 11:19:26PM +0300, Ilya Verbin wrote:
> > '#pragma omp critical (name)' can be placed in the function, marked
> > with '#pragma omp declare target', in this case the corresponding node
> > should be marked as offloadable too.
> > Bootstrapped/regtested on x86_64-linux and i686-linux, ok for trunk?
> 
> Please add a testcase for this.

Here is the testcase, based on critical-2.c, ok for trunk?


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.



  -- Ilya

Comments

Jakub Jelinek Nov. 28, 2014, 8:52 a.m. UTC | #1 
On Fri, Nov 28, 2014 at 02:52:11AM +0300, Ilya Verbin wrote:
> On 21 Nov 21:36, Jakub Jelinek wrote:
> > On Fri, Nov 21, 2014 at 11:19:26PM +0300, Ilya Verbin wrote:
> > > '#pragma omp critical (name)' can be placed in the function, marked
> > > with '#pragma omp declare target', in this case the corresponding node
> > > should be marked as offloadable too.
> > > Bootstrapped/regtested on x86_64-linux and i686-linux, ok for trunk?
> > 
> > Please add a testcase for this.
> 
> Here is the testcase, based on critical-2.c, ok for trunk?
> 
> 
> 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.

Ok, thanks.

	Jakub
diff mbox

Patch

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index dff1504..621958c 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9387,16 +9387,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);
@@ -9404,6 +9394,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/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;
+}