Hi Thomas!
This patch introduces support of OpenACC loop directive (and combined
directives) in C front-end up to GENERIC. Currently no clause is allowed.
This patch is necessary to finish implementation of OpenACC 1.0 in
fortran front-end. As you know, OpenACC fortran implementation does
parsing and resolving of loop directive but doesn't transformation to
GENERIC.
Bootstraped and tested with no new regressions on x86_64-unknown-linux-gnu.
OK for gomp4 branch?
--
Ilmir.
>From 1adbb9d2e504c7acac8be89b053fa677cf285b42 Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usma...@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
