From 1adbb9d2e504c7acac8be89b053fa677cf285b42 Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usmanov@samsung.com>
Date: Tue, 18 Mar 2014 16:13:14 +0400
Subject: [PATCH] Initial support of OpenACC loop

---
	Initial support of OpenACC loop directive in C front-end.

	gcc/
	* tree.def (OACC_LOOP): New tree code.
	* tree-pretty-print.c (dump_generic_node): Show it.
	* tree.h (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): New macros.
	* c-family/c-pragma.c (oacc_pragmas): Add loop directive.
	* c-family/c-pragma.h (enum pragma_kind): Likewise.
	* c/c-parser.c (c_parser_oacc_loop): New function.
	(c_parser_oacc_kernels): Parse combined directive.
	(c_parser_oacc_parallel): Likewise.
	(c_parser_omp_construct): Parse loop directive.
	* doc/generic.texi: Document loop directive.
	* gimplify.c (is_gimple_stmt, gimplify_expr): Stub gimplification of 
	loop directive and combined directives.
	* testsuite/c-c++-common/goacc/loop-1.c: New test.

diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index c8baba4..f99b087 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1172,6 +1172,7 @@ static const struct omp_pragma_def oacc_pragmas[] = {
   { "data", PRAGMA_OACC_DATA },
   { "kernels", PRAGMA_OACC_KERNELS },
   { "parallel", PRAGMA_OACC_PARALLEL },
+  { "loop", PRAGMA_OACC_LOOP },
 };
 static const struct omp_pragma_def omp_pragmas[] = {
   { "atomic", PRAGMA_OMP_ATOMIC },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index d55a511..16c74a9 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -30,6 +30,7 @@ typedef enum pragma_kind {
   PRAGMA_OACC_DATA,
   PRAGMA_OACC_KERNELS,
   PRAGMA_OACC_PARALLEL,
+  PRAGMA_OACC_LOOP,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
   PRAGMA_OMP_CANCEL,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 3d8c0de..2300c6c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11404,6 +11404,8 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
   return stmt;
 }
 
+static tree c_parser_oacc_loop (location_t, c_parser *, char *);
+
 /* OpenACC 2.0:
    # pragma acc kernels oacc-kernels-clause[optseq] new-line
      structured-block
@@ -11424,12 +11426,28 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
 
 static tree
-c_parser_oacc_kernels (location_t loc, c_parser *parser)
+c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
 {
-  tree stmt, clauses, block;
+  tree stmt, clauses = NULL_TREE, block;
+
+  strcat (p_name, " kernels");
+
+  if (c_parser_next_token_is (parser, CPP_NAME))
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      if (strcmp (p, "loop") == 0)
+	{
+	  c_parser_consume_token (parser);
+	  block = c_begin_omp_parallel ();
+	  c_parser_oacc_loop (loc, parser, p_name);
+	  stmt = c_finish_oacc_kernels (loc, clauses, block);
+	  OACC_KERNELS_COMBINED (stmt) = 1;
+	  return stmt;
+	}
+    }
 
   clauses =  c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
-					"#pragma acc kernels");
+					p_name);
 
   block = c_begin_omp_parallel ();
   add_stmt (c_parser_omp_structured_block (parser));
@@ -11459,12 +11477,28 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
 
 static tree
-c_parser_oacc_parallel (location_t loc, c_parser *parser)
+c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
 {
-  tree stmt, clauses, block;
+  tree stmt, clauses = NULL_TREE, block;
+
+  strcat (p_name, " parallel");
+
+  if (c_parser_next_token_is (parser, CPP_NAME))
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      if (strcmp (p, "loop") == 0)
+	{
+	  c_parser_consume_token (parser);
+	  block = c_begin_omp_parallel ();
+	  c_parser_oacc_loop (loc, parser, p_name);
+	  stmt = c_finish_oacc_parallel (loc, clauses, block);
+	  OACC_PARALLEL_COMBINED (stmt) = 1;
+	  return stmt;
+	}
+    }
 
   clauses =  c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
-					"#pragma acc parallel");
+					p_name);
 
   block = c_begin_omp_parallel ();
   add_stmt (c_parser_omp_structured_block (parser));
@@ -12243,6 +12277,32 @@ omp_split_clauses (location_t loc, enum tree_code code,
       cclauses[i] = c_finish_omp_clauses (cclauses[i]);
 }
 
+/* OpenACC 2.0:
+   # pragma acc loop oacc-loop-clause[optseq] new-line
+     structured-block
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_LOOP_CLAUSE_MASK	PRAGMA_OMP_CLAUSE_NONE
+
+static tree
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+{
+  tree block, clauses, ret;
+
+  strcat (p_name, " loop");
+
+  clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+
+  block = c_begin_compound_stmt (true);
+  ret = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+  block = c_end_compound_stmt (loc, block, true);
+  add_stmt (block);
+
+  return ret;
+}
+
 /* OpenMP 4.0:
    #pragma omp simd simd-clause[optseq] new-line
      for-loop
@@ -13757,10 +13817,16 @@ c_parser_omp_construct (c_parser *parser)
       stmt = c_parser_oacc_data (loc, parser);
       break;
     case PRAGMA_OACC_KERNELS:
-      stmt = c_parser_oacc_kernels (loc, parser);
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_kernels (loc, parser, p_name);
       break;
     case PRAGMA_OACC_PARALLEL:
-      stmt = c_parser_oacc_parallel (loc, parser);
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_parallel (loc, parser, p_name);
+      break;
+    case PRAGMA_OACC_LOOP:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_loop (loc, parser, p_name);
       break;
     case PRAGMA_OMP_ATOMIC:
       c_parser_omp_atomic (loc, parser);
diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi
index ce14620..0a77a86 100644
--- a/gcc/doc/generic.texi
+++ b/gcc/doc/generic.texi
@@ -2054,6 +2054,7 @@ edge.  Rethrowing the exception is represented using @code{RESX_EXPR}.
 @tindex OACC_PARALLEL
 @tindex OACC_KERNELS
 @tindex OACC_DATA
+@tindex OACC_LOOP
 @tindex OACC_HOST_DATA
 @tindex OACC_DECLARE
 @tindex OACC_UPDATE
@@ -2090,6 +2091,10 @@ Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
 
 Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
 
+@item OACC_LOOP
+
+Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}.
+
 @item OACC_HOST_DATA
 
 Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index f3c34f9..3992737 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -4363,6 +4363,7 @@ is_gimple_stmt (tree t)
     case OMP_FOR:
     case OMP_SIMD:
     case CILK_SIMD:
+    case OACC_LOOP:
     case OMP_DISTRIBUTE:
     case OMP_SECTIONS:
     case OMP_SECTION:
@@ -8047,6 +8048,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case OACC_EXIT_DATA:
 	case OACC_WAIT:
 	case OACC_CACHE:
+	case OACC_LOOP:
 	  sorry ("directive not yet implemented");
 	  ret = GS_ALL_DONE;
 	  break;
@@ -8068,9 +8070,23 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = gimplify_omp_for (expr_p, pre_p);
 	  break;
 
-	case OACC_DATA:
 	case OACC_KERNELS:
+	  if (OACC_KERNELS_COMBINED (*expr_p))
+	    sorry ("directive not yet implemented");
+	  else
+	    gimplify_omp_workshare (expr_p, pre_p);
+	  ret = GS_ALL_DONE;
+	  break;
+
 	case OACC_PARALLEL:
+	  if (OACC_PARALLEL_COMBINED (*expr_p))
+	    sorry ("directive not yet implemented");
+	  else
+	    gimplify_omp_workshare (expr_p, pre_p);
+	  ret = GS_ALL_DONE;
+	  break;
+
+	case OACC_DATA:
 	case OMP_SECTIONS:
 	case OMP_SINGLE:
 	case OMP_TARGET:
@@ -8473,6 +8489,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 		  && code != OACC_CACHE
 		  && code != OMP_CRITICAL
 		  && code != OMP_FOR
+		  && code != OACC_LOOP
 		  && code != OMP_MASTER
 		  && code != OMP_TASKGROUP
 		  && code != OMP_ORDERED
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
new file mode 100644
index 0000000..bb36252
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -0,0 +1,89 @@
+/* { dg-do compile } */
+
+int test1()
+{
+  int i, j, k, b[10];
+  int a[30];
+  double d;
+  float r;
+  i = 0;
+  #pragma acc loop
+  while(1)  /* { dg-error "for statement expected" } */
+    {
+      if (i > 0) break; 
+      i = i + 1;
+    }
+  i = 0;
+  #pragma acc loop
+  for(;;)  /* { dg-error "expected iteration declaration or initialization" } */
+    {
+      if (i > 0) break; /* { dg-error "break statement used" } */
+      i = i + 1;
+    }
+  i = 0;
+  #pragma acc loop
+  do  /* { dg-error "for statement expected" } */
+    {
+      i = i + 1;
+    }
+  while (i < 4);
+  #pragma acc loop
+  while (i < 8)  /* { dg-error "for statement expected" } */
+    {
+      i = i + 1;
+    }
+  #pragma acc loop
+  for (d = 1; d < 30; d+= 6)  /* { dg-error "invalid type for iteration variable" } */
+    {
+      i = d;
+      a[i] = 1;
+    }
+  #pragma acc loop
+  for (d = 1; d < 30; d+= 5)  /* { dg-error "invalid type for iteration variable" } */
+    {
+      i = d;
+      a[i] = 2;
+    }
+  #pragma acc loop
+  for (i = 1; i < 30; i++ )
+    if (i == 16) break; /* { dg-error "break statement used" } */
+
+  #pragma acc loop
+  for(i = 1; i < 30; i++)
+    {
+      for(j = 5; j < 10; j++)
+        {
+          /* TODO: there must be error.  */
+          if (i == 6 && j == 7) goto outer; 
+        }
+    }
+outer:
+
+/* different types of for loop are allowed */
+  #pragma acc loop
+  for (i = 1; i < 10; i++)
+    {
+    }
+  #pragma acc loop
+  for (i = 1; i < 10; i+=2)
+    {
+      a[i] = i;
+    }
+
+  /* after loop directive must be loop */
+  #pragma acc loop
+    a[1] = 1; /* { dg-error "for statement expected" } */
+    for (i = 1; i < 10; i++)
+      ;
+  /* combined directives may be used*/
+  #pragma acc parallel loop
+  for(i = 1; i < 10; i++)
+    {
+    }
+  #pragma acc kernels loop
+  for(i = 1; i < 10; i++)
+    {
+    }
+  return 0;
+}
+/* { dg-excess-errors "directive not yet implemented" } */
\ No newline at end of file
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index d1a3300..49e5f6c 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -2538,6 +2538,10 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
       pp_string (buffer, "#pragma simd");
       goto dump_omp_loop;
 
+    case OACC_LOOP:
+      pp_string (buffer, "#pragma acc loop");
+      goto dump_omp_loop;
+
     case OMP_DISTRIBUTE:
       pp_string (buffer, "#pragma omp distribute");
       goto dump_omp_loop;
diff --git a/gcc/tree.def b/gcc/tree.def
index 623ebb0..d9e4eb41 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1076,6 +1076,10 @@ DEFTREECODE (OMP_SIMD, "omp_simd", tcc_statement, 6)
    Operands like for OMP_FOR.  */
 DEFTREECODE (CILK_SIMD, "cilk_simd", tcc_statement, 6)
 
+/* OpenACC - #pragma acc loop [clause1 ... clauseN]
+   Operands like for OMP_FOR.  */
+DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6)
+
 /* OpenMP - #pragma omp distribute [clause1 ... clauseN]
    Operands like for OMP_FOR.  */
 DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
diff --git a/gcc/tree.h b/gcc/tree.h
index d98165c..6668895 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1269,6 +1269,15 @@ extern void protected_set_expr_location (tree, location_t);
 #define OMP_SECTION_LAST(NODE) \
   (OMP_SECTION_CHECK (NODE)->base.private_flag)
 
+/* True on an OACC_KERNELS statement if is represents combined kernels loop
+   directive.  */
+#define OACC_KERNELS_COMBINED(NODE) \
+  (OACC_KERNELS_CHECK (NODE)->base.private_flag)
+
+/* Like OACC_KERNELS_COMBINED, but for parallel loop directive.  */
+#define OACC_PARALLEL_COMBINED(NODE) \
+  (OACC_PARALLEL_CHECK (NODE)->base.private_flag)
+
 /* True on an OMP_PARALLEL statement if it represents an explicit
    combined parallel work-sharing constructs.  */
 #define OMP_PARALLEL_COMBINED(NODE) \
-- 
1.8.3.2

