This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [OpenACC 4/11] C FE changes


On 10/23/2015 02:31 PM, Cesar Philippidis wrote:
> On 10/23/2015 01:31 PM, Jakub Jelinek wrote:
>> On Fri, Oct 23, 2015 at 01:17:07PM -0700, Cesar Philippidis wrote:
>>> Good idea, thanks. This patch also corrects the problems parsing weird
>>> combinations of num, static and length arguments that you mentioned
>>> elsewhere.
>>>
>>> Is this OK for trunk?
>>
>> I'd strongly prefer to see always patches accompanied by testcases.
>>
>>> +	  loc = c_parser_peek_token (parser)->location;
>>> +	  op_to_parse = &op0;
>>> +
>>> +	  if ((c_parser_next_token_is (parser, CPP_NAME)
>>> +	       || c_parser_next_token_is (parser, CPP_KEYWORD))
>>> +	      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
>>> +	    {
>>> +	      tree name_kind = c_parser_peek_token (parser)->value;
>>> +	      const char *p = IDENTIFIER_POINTER (name_kind);
>>
>> I think I'd prefer not to peek at this at all if it is RID_STATIC,
>> so perhaps just have (and name_kind is weird):
>> 	      else
>> 		{
>> 		  tree val = c_parser_peek_token (parser)->value;
>> 		  if (strcmp (id, IDENTIFIER_POINTER (val)) == 0)
>> 		    {
>> 		      c_parser_consume_token (parser);  /* id  */
>> 		      c_parser_consume_token (parser);  /* ':'  */
>> 		    }
>> 		  else
>> 		    {
>> ...
>> 		    }
>> 		}
>> ?
> 
> My plan over here was try and catch any arguments with a colon. But that
> fell threw because...
> 
>>> +	      if (kind == OMP_CLAUSE_GANG
>>> +		  && c_parser_next_token_is_keyword (parser, RID_STATIC))
>>> +		{
>>> +		  c_parser_consume_token (parser); /* static  */
>>> +		  c_parser_consume_token (parser); /* ':'  */
>>> +
>>> +		  op_to_parse = &op1;
>>> +		  if (c_parser_next_token_is (parser, CPP_MULT))
>>> +		    {
>>> +		      c_parser_consume_token (parser);
>>> +		      *op_to_parse = integer_minus_one_node;
>>> +
>>> +		      /* Consume a comma if present.  */
>>> +		      if (c_parser_next_token_is (parser, CPP_COMMA))
>>> +			c_parser_consume_token (parser);
>>
>> Doesn't this mean that you happily parse
>> gang (static: * abc)
>> or
>> gang (static:*num:1)
>> etc.?  I'd say the comma should be non-optional (i.e. either accept
>> CPP_COMMA, or CPP_CLOSE_PARENT, but nothing else) in that case (at least,
>> when in OpenMP grammar something is *-list it is meant to be comma
>> separated).
> 
> I'm not handling commas properly. My next patch is going to handle the
> static argument separately.
> 
>>> +	  /* Consume a comma if present.  */
>>> +	  if (c_parser_next_token_is (parser, CPP_COMMA))
>>> +	    c_parser_consume_token (parser);
>>
>> Similarly this means
>> gang (num: 5 static: *)
>> is accepted.  If it is valid, then again it should have testsuite coverage.
> 
> I'll include a test case for this with the next patch.

Here's the updated patch. Hopefully I addressed everything. Thank you
for suggesting all of those test cases.

Is this OK for trunk?

Cesar

2015-10-23  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>

	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.

2015-10-23  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/testsuite/
	* c-c++-common/goacc/loop-shape.c (int main):

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index c8c6a2d..7d2baa9 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11188,6 +11188,156 @@ 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 op0 = NULL_TREE, op1 = 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))
+    {
+      tree *op_to_parse = &op0;
+      c_token *next;
+
+      c_parser_consume_token (parser);
+
+      do
+	{
+	  next = c_parser_peek_token (parser);
+	  loc = next->location;
+	  op_to_parse = &op0;
+
+	  /* Consume a comma if present.  */
+	  if (c_parser_next_token_is (parser, CPP_COMMA))
+	    {
+	      if (op0 == NULL && op1 == NULL)
+		{
+		  c_parser_error (parser, "unexpected argument");
+		  goto cleanup_error;
+		}
+
+	      c_parser_consume_token (parser);
+	    }
+
+	  /* First extract any length:, num: and static: arguments.  */
+
+	  /* Gang static argument.  */
+	  if (c_parser_next_token_is (parser, CPP_KEYWORD)
+	      && c_parser_next_token_is_keyword (parser, RID_STATIC))
+	    {
+	      if (kind != OMP_CLAUSE_GANG)
+		{
+		  c_parser_error (parser, "invalid %<static%> argument");
+		  goto cleanup_error;
+		}
+
+	      c_parser_consume_token (parser);
+
+	      if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+		goto cleanup_error;
+
+	      op_to_parse = &op1;
+	      if (*op_to_parse != NULL)
+		{
+		  c_parser_error (parser, "too many %<static%> arguements");
+		  goto cleanup_error;
+		}
+
+	      /* Check for the '*' argument.  */
+	      if (c_parser_next_token_is (parser, CPP_MULT))
+		{
+		  c_parser_consume_token (parser);
+		  *op_to_parse = integer_minus_one_node;
+		  continue;
+		}
+	    }
+	  /* 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 (*op_to_parse != NULL_TREE)
+	    {
+	      c_parser_error (parser, "unexpected argument");
+	      goto cleanup_error;
+	    }
+
+	  tree expr = c_parser_expr_no_commas (parser, NULL).value;
+	  if (expr == error_mark_node)
+	    goto cleanup_error;
+
+	  mark_exp_read (expr);
+	  *op_to_parse = expr;
+	}
+      while (c_parser_next_token_is (parser, CPP_COMMA)
+	     && !c_parser_next_token_is (parser, CPP_CLOSE_PAREN));
+
+      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 (op0)
+    OMP_CLAUSE_OPERAND (c, 0) = op0;
+  if (op1)
+    OMP_CLAUSE_OPERAND (c, 1) = op1;
+  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 ATTRIBUTE_UNUSED,
+			     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
@@ -12393,6 +12543,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";
@@ -12429,6 +12584,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";
@@ -12477,6 +12637,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";
@@ -12485,6 +12655,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;
@@ -13015,6 +13190,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
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
new file mode 100644
index 0000000..1053361
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -0,0 +1,218 @@
+/* Exercise *_parser_oacc_shape_clause by checking various combinations
+   of gang, worker and vector clause arguments.  */
+
+/* { dg-compile } */
+
+int main ()
+{
+  int i;
+  int v, w;
+  int length, num;
+
+  /* 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 "unexpected argument" } */
+  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 "unexpected argument" } */
+  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++)
+    ;
+
+  
+  return 0;
+}

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]