diff mbox

[GOMP4,COMMITTED] OpenACC if clause.

Message ID 1402570722-17540-1-git-send-email-thomas@codesourcery.com
State New
Headers show

Commit Message

Thomas Schwinge June 12, 2014, 10:58 a.m. UTC 
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>

	gcc/c/
	* c-parser.c (c_parser_oacc_all_clauses): Handle
	PRAGMA_OMP_CLAUSE_IF.
	(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
	(OACC_PARALLEL_CLAUSE_MASK, OACC_UPDATE_CLAUSE_MASK): Add it.
	gcc/
	* omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_IF.
	(expand_oacc_offload, expand_omp_target): Handle it.
	gcc/testsuite/
	* c-c++-common/goacc/if-clause-1.c: New file.
	* c-c++-common/goacc/if-clause-2.c: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211510 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                             |  6 ++
 gcc/c/ChangeLog.gomp                           |  8 +++
 gcc/c/c-parser.c                               | 10 +++-
 gcc/omp-low.c                                  | 81 +++++++++++++++++++++-----
 gcc/testsuite/ChangeLog.gomp                   |  5 ++
 gcc/testsuite/c-c++-common/goacc/if-clause-1.c |  8 +++
 gcc/testsuite/c-c++-common/goacc/if-clause-2.c | 11 ++++
 7 files changed, 112 insertions(+), 17 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/if-clause-1.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/if-clause-2.c
diff mbox

Patch

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index be1aa16..2abe179 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@ 
+2014-06-12  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_IF.
+	(expand_oacc_offload, expand_omp_target): Handle it.
+
 2014-06-06  Thomas Schwinge  <thomas@codesourcery.com>
 	    James Norris  <jnorris@codesourcery.com>
 
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index f1e45f3..108ce3e 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,11 @@ 
+2014-06-12  Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* c-parser.c (c_parser_oacc_all_clauses): Handle
+	PRAGMA_OMP_CLAUSE_IF.
+	(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+	(OACC_PARALLEL_CLAUSE_MASK, OACC_UPDATE_CLAUSE_MASK): Add it.
+
 2014-06-06  Thomas Schwinge  <thomas@codesourcery.com>
 	    James Norris  <jnorris@codesourcery.com>
 
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index bf4bad62..6269923 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10203,7 +10203,7 @@  c_parser_omp_clause_final (c_parser *parser, tree list)
   return list;
 }
 
-/* OpenMP 2.5:
+/* OpenACC, OpenMP 2.5:
    if ( expression ) */
 
 static tree
@@ -11295,6 +11295,10 @@  c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "host";
 	  break;
+	case PRAGMA_OMP_CLAUSE_IF:
+	  clauses = c_parser_omp_clause_if (parser, clauses);
+	  c_name = "if";
+	  break;
 	case PRAGMA_OMP_CLAUSE_NUM_GANGS:
 	  clauses = c_parser_omp_clause_num_gangs (parser, clauses);
 	  c_name = "num_gangs";
@@ -11614,6 +11618,7 @@  c_parser_omp_structured_block (c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN)	\
@@ -11649,6 +11654,7 @@  c_parser_oacc_data (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN)	\
@@ -11727,6 +11733,7 @@  c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_GANGS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT)		\
@@ -11775,6 +11782,7 @@  c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
 #define OACC_UPDATE_CLAUSE_MASK						\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) )
 
 static void
diff --git gcc/omp-low.c gcc/omp-low.c
index 6fd0f30..958f116 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1634,12 +1634,6 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
 	  break;
 
-	case OMP_CLAUSE_IF:
-	  if (is_gimple_omp_oacc_specifically (ctx->stmt))
-	    {
-	      sorry ("clause not supported yet");
-	      break;
-	    }
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
@@ -1649,6 +1643,7 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEPEND:
 	  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
@@ -1916,12 +1911,6 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    }
 	  break;
 
-	case OMP_CLAUSE_IF:
-	  if (is_gimple_omp_oacc_specifically (ctx->stmt))
-	    {
-	      sorry ("clause not supported yet");
-	      break;
-	    }
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_COPYIN:
 	case OMP_CLAUSE_DEFAULT:
@@ -1945,6 +1934,7 @@  scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
 	  gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
@@ -5115,7 +5105,7 @@  expand_oacc_offload (struct omp_region *region)
   /* Emit a library call to launch CHILD_FN.  */
   tree t1, t2, t3, t4,
     t_num_gangs, t_num_workers, t_vector_length,
-    device, c, clauses;
+    device, cond, c, clauses;
   enum built_in_function start_ix;
   location_t clause_loc;
   tree (*gimple_omp_clauses) (const_gimple);
@@ -5165,9 +5155,15 @@  expand_oacc_offload (struct omp_region *region)
       break;
     }
 
-  /* By default, the value of DEVICE is -1 (let runtime library choose).  */
+  /* By default, the value of DEVICE is -1 (let runtime library choose)
+     and there is no conditional.  */
+  cond = NULL_TREE;
   device = build_int_cst (integer_type_node, -1);
 
+  c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+  if (c)
+    cond = OMP_CLAUSE_IF_EXPR (c);
+
   c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
   gcc_assert (c == NULL);
   if (c)
@@ -5181,6 +5177,61 @@  expand_oacc_offload (struct omp_region *region)
   /* Ensure 'device' is of the correct type.  */
   device = fold_convert_loc (clause_loc, integer_type_node, device);
 
+  /* If we found the clause 'if (cond)', build
+     (cond ? device : -2).  */
+  if (cond)
+    {
+      cond = gimple_boolify (cond);
+
+      basic_block cond_bb, then_bb, else_bb;
+      edge e;
+      tree tmp_var;
+
+      tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
+      /* Preserve indentation of expand_omp_target.  */
+      if (0)
+	{
+	  gsi = gsi_last_bb (new_bb);
+	  gsi_prev (&gsi);
+	  e = split_block (new_bb, gsi_stmt (gsi));
+	}
+      else
+	e = split_block (new_bb, NULL);
+      cond_bb = e->src;
+      new_bb = e->dest;
+      remove_edge (e);
+
+      then_bb = create_empty_bb (cond_bb);
+      else_bb = create_empty_bb (then_bb);
+      set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+      set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
+
+      stmt = gimple_build_cond_empty (cond);
+      gsi = gsi_last_bb (cond_bb);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      gsi = gsi_start_bb (then_bb);
+      stmt = gimple_build_assign (tmp_var, device);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      gsi = gsi_start_bb (else_bb);
+      stmt = gimple_build_assign (tmp_var,
+				  build_int_cst (integer_type_node, -2));
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+      make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
+      if (current_loops)
+	{
+	  add_bb_to_loop (then_bb, cond_bb->loop_father);
+	  add_bb_to_loop (else_bb, cond_bb->loop_father);
+	}
+      make_edge (then_bb, new_bb, EDGE_FALLTHRU);
+      make_edge (else_bb, new_bb, EDGE_FALLTHRU);
+
+      device = tmp_var;
+    }
+
   gsi = gsi_last_bb (new_bb);
   t = gimple_omp_data_arg (entry_stmt);
   if (t == NULL)
@@ -8661,7 +8712,6 @@  expand_omp_target (struct omp_region *region)
   device = build_int_cst (integer_type_node, -1);
 
   c = find_omp_clause (clauses, OMP_CLAUSE_IF);
-  gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
   if (c)
     cond = OMP_CLAUSE_IF_EXPR (c);
 
@@ -8684,7 +8734,6 @@  expand_omp_target (struct omp_region *region)
      (cond ? device : -2).  */
   if (cond)
     {
-      gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
       cond = gimple_boolify (cond);
 
       basic_block cond_bb, then_bb, else_bb;
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 17493ce..9b65cc9 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@ 
+2014-06-12  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/if-clause-1.c: New file.
+	* c-c++-common/goacc/if-clause-2.c: Likewise.
+
 2014-06-06  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-c++-common/goacc/pragma_context.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc/if-clause-1.c gcc/testsuite/c-c++-common/goacc/if-clause-1.c
new file mode 100644
index 0000000..c078596
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/if-clause-1.c
@@ -0,0 +1,8 @@ 
+void
+f (void)
+{
+  struct { int i; } *p;
+#pragma acc data copyout(p) if(1) if(1) /* { dg-error "too many 'if' clauses" } */
+  ;
+#pragma acc update device(p) if(*p) /* { dg-error "used struct type value where scalar is required" } */
+}
diff --git gcc/testsuite/c-c++-common/goacc/if-clause-2.c gcc/testsuite/c-c++-common/goacc/if-clause-2.c
new file mode 100644
index 0000000..5ab8459
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/if-clause-2.c
@@ -0,0 +1,11 @@ 
+void
+f (short c)
+{
+#pragma acc parallel if(c)
+  ;
+#pragma acc kernels if(c)
+  ;
+#pragma acc data if(c)
+  ;
+#pragma acc update device(c) if(c)
+}