c-parser.c (c_parser_oacc_shape_clause): New.

	gcc/c/
	* c-parser.c (c_parser_oacc_shape_clause): New.
	(c_parser_oacc_simple_clause): New.
	(c_parser_oacc_all_clauses): Add auto, gang, seq, vector, worker.
	(OACC_LOOP_CLAUSE_MASK): Add gang, worker, vector, auto, seq.

	gcc/testsuite/
	* c-c++-common/goacc/loop-shape.c: New test.

Co-Authored-By: Bernd Schmidt <bernds@redhat.com>
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Joseph Myers <joseph@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r229461

gcc/c/ChangeLog

+2015-10-27  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+	    Joseph Myers  <joseph@codesourcery.com>
+	    Julian Brown  <julian@codesourcery.com>
+	    Bernd Schmidt  <bschmidt@redhat.com>
+
+	* c-parser.c (c_parser_oacc_shape_clause): New.
+	(c_parser_oacc_simple_clause): New.
+	(c_parser_oacc_all_clauses): Add auto, gang, seq, vector, worker.
+	(OACC_LOOP_CLAUSE_MASK): Add gang, worker, vector, auto, seq.
+
 2015-10-27  Thomas Schwinge  <thomas@codesourcery.com>
 	    James Norris  <jnorris@codesourcery.com>
 	    Cesar Philippidis  <cesar@codesourcery.com>

gcc/c/c-parser.c

@@ -11187,6 +11187,167 @@ c_parser_omp_clause_num_workers (c_parser *parser, tree list)
 }
 
 /* OpenACC:
+
+    gang [( gang-arg-list )]
+    worker [( [num:] int-expr )]
+    vector [( [length:] int-expr )]
+
+  where gang-arg is one of:
+
+    [num:] int-expr
+    static: size-expr
+
+  and size-expr may be:
+
+    *
+    int-expr
+*/
+
+static tree
+c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
+			    const char *str, tree list)
+{
+  const char *id = "num";
+  tree ops[2] = { NULL_TREE, NULL_TREE }, c;
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  if (kind == OMP_CLAUSE_VECTOR)
+    id = "length";
+
+  if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+    {
+      c_parser_consume_token (parser);
+
+      do
+	{
+	  c_token *next = c_parser_peek_token (parser);
+	  int idx = 0;
+
+	  /* Gang static argument.  */
+	  if (kind == OMP_CLAUSE_GANG
+	      && c_parser_next_token_is_keyword (parser, RID_STATIC))
+	    {
+	      c_parser_consume_token (parser);
+
+	      if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+		goto cleanup_error;
+
+	      idx = 1;
+	      if (ops[idx] != NULL_TREE)
+		{
+		  c_parser_error (parser, "too many %<static%> arguments");
+		  goto cleanup_error;
+		}
+
+	      /* Check for the '*' argument.  */
+	      if (c_parser_next_token_is (parser, CPP_MULT))
+		{
+		  c_parser_consume_token (parser);
+		  ops[idx] = integer_minus_one_node;
+
+		  if (c_parser_next_token_is (parser, CPP_COMMA))
+		    {
+		      c_parser_consume_token (parser);
+		      continue;
+		    }
+		  else
+		    break;
+		}
+	    }
+	  /* Worker num: argument and vector length: arguments.  */
+	  else if (c_parser_next_token_is (parser, CPP_NAME)
+		   && strcmp (id, IDENTIFIER_POINTER (next->value)) == 0
+		   && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+	    {
+	      c_parser_consume_token (parser);  /* id  */
+	      c_parser_consume_token (parser);  /* ':'  */
+	    }
+
+	  /* Now collect the actual argument.  */
+	  if (ops[idx] != NULL_TREE)
+	    {
+	      c_parser_error (parser, "unexpected argument");
+	      goto cleanup_error;
+	    }
+
+	  location_t expr_loc = c_parser_peek_token (parser)->location;
+	  tree expr = c_parser_expr_no_commas (parser, NULL).value;
+	  if (expr == error_mark_node)
+	    goto cleanup_error;
+
+	  mark_exp_read (expr);
+	  expr = c_fully_fold (expr, false, NULL);
+
+	  /* Attempt to statically determine when the number isn't a
+	     positive integer.  */
+
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)))
+	    {
+	      c_parser_error (parser, "expected integer expression");
+	      return list;
+	    }
+
+	  tree c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, expr,
+				    build_int_cst (TREE_TYPE (expr), 0));
+	  if (c == boolean_true_node)
+	    {
+	      warning_at (loc, 0,
+			  "%<%s%> value must be positive", str);
+	      expr = integer_one_node;
+	    }
+
+	  ops[idx] = expr;
+
+	  if (kind == OMP_CLAUSE_GANG
+	      && c_parser_next_token_is (parser, CPP_COMMA))
+	    {
+	      c_parser_consume_token (parser);
+	      continue;
+	    }
+	  break;
+	}
+      while (1);
+
+      if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+	goto cleanup_error;
+    }
+
+  check_no_duplicate_clause (list, kind, str);
+
+  c = build_omp_clause (loc, kind);
+
+  if (ops[1])
+    OMP_CLAUSE_OPERAND (c, 1) = ops[1];
+
+  OMP_CLAUSE_OPERAND (c, 0) = ops[0];
+  OMP_CLAUSE_CHAIN (c) = list;
+
+  return c;
+
+ cleanup_error:
+  c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+  return list;
+}
+
+/* OpenACC:
+   auto
+   independent
+   nohost
+   seq */
+
+static tree
+c_parser_oacc_simple_clause (c_parser *parser, enum omp_clause_code code,
+			     tree list)
+{
+  check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
+
+  tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
+  OMP_CLAUSE_CHAIN (c) = list;
+
+  return c;
+}
+
+/* OpenACC:
    async [( int-expr )] */
 
 static tree
@@ -12392,6 +12553,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_clause_async (parser, clauses);
 	  c_name = "async";
 	  break;
+	case PRAGMA_OACC_CLAUSE_AUTO:
+	  clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO,
+						clauses);
+	  c_name = "auto";
+	  break;
 	case PRAGMA_OACC_CLAUSE_COLLAPSE:
 	  clauses = c_parser_omp_clause_collapse (parser, clauses);
 	  c_name = "collapse";
@@ -12428,6 +12594,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_firstprivate (parser, clauses);
 	  c_name = "firstprivate";
 	  break;
+	case PRAGMA_OACC_CLAUSE_GANG:
+	  c_name = "gang";
+	  clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+						c_name, clauses);
+	  break;
 	case PRAGMA_OACC_CLAUSE_HOST:
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "host";
@@ -12476,6 +12647,16 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "self";
 	  break;
+	case PRAGMA_OACC_CLAUSE_SEQ:
+	  clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
+						clauses);
+	  c_name = "seq";
+	  break;
+	case PRAGMA_OACC_CLAUSE_VECTOR:
+	  c_name = "vector";
+	  clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+						c_name,	clauses);
+	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
 	  clauses = c_parser_omp_clause_vector_length (parser, clauses);
 	  c_name = "vector_length";
@@ -12484,6 +12665,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_clause_wait (parser, clauses);
 	  c_name = "wait";
 	  break;
+	case PRAGMA_OACC_CLAUSE_WORKER:
+	  c_name = "worker";
+	  clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+						c_name, clauses);
+	  break;
 	default:
 	  c_parser_error (parser, "expected %<#pragma acc%> clause");
 	  goto saw_error;
@@ -12960,6 +13146,11 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
 
 #define OACC_LOOP_CLAUSE_MASK						\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
 
 static tree

gcc/testsuite/ChangeLog

+2015-10-27  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* c-c++-common/goacc/loop-shape.c: New test.
+
 2015-10-27  Nathan Sidwell <nathan@codesourcery.com>
 	    Cesar Philippidis  <cesar@codesourcery.com>
 

gcc/testsuite/c-c++-common/goacc/loop-shape.c

+/* Exercise *_parser_oacc_shape_clause by checking various combinations
+   of gang, worker and vector clause arguments.  */
+
+/* { dg-compile } */
+
+int main ()
+{
+  int i;
+  int v = 32, w = 19;
+  int length = 1, num = 5;
+
+  /* Valid uses.  */
+
+  #pragma acc kernels
+  #pragma acc loop gang worker vector
+  for (i = 0; i < 10; i++)
+    ;
+  
+  #pragma acc kernels
+  #pragma acc loop gang(26)
+  for (i = 0; i < 10; i++)
+    ;
+  
+  #pragma acc kernels
+  #pragma acc loop gang(v)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length: 16)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length: v)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(16)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(v)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num: 16)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num: v)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(16)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(v)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static: 16, num: 5)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static: v, num: w)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num, static: 6)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static: 5, num)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(1, static:*)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static:*, 1)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(1, static:*)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num: 5, static: 4)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num: v, static: w)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num, static:num)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length:length)
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num:length)
+  for (i = 0; i < 10; i++)
+    ;  
+
+  #pragma acc kernels
+  #pragma acc loop worker(num:num)
+  for (i = 0; i < 10; i++)
+    ;  
+
+  /* Invalid uses.  */
+  
+  #pragma acc kernels
+  #pragma acc loop gang(16, 24) /* { dg-error "unexpected argument" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(v, w) /* { dg-error "unexpected argument" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num: 1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num: 1 num:2, num:3, 4) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num, num:5) /* { dg-error "unexpected argument" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(length:num) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(5, length:length) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(num:length) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(length:5) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(1, num:2) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num: 5 static: *) /* { dg-error "expected '.' before" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(,static: *) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(,length:5) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(,num:10) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(,10) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(,10) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(,10) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(-12) /* { dg-warning "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num:-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(num:1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static:-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop gang(static:1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num:-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop worker(num:1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length:-1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  #pragma acc kernels
+  #pragma acc loop vector(length:1.0) /* { dg-error "" } */
+  for (i = 0; i < 10; i++)
+    ;
+
+  return 0;
+}