This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[C++ PATCH] [gomp4] Initial OpenACC support to C++ front-end
- From: Ilmir Usmanov <i dot usmanov at samsung dot com>
- To: Thomas Schwinge <thomas at codesourcery dot com>
- Cc: Evgeny Gavrin <e dot gavrin at samsung dot com>, Slava Garbuzov <v dot garbuzov at samsung dot com>, gcc-patches at gcc dot gnu dot org
- Date: Thu, 13 Mar 2014 21:05:17 +0400
- Subject: [C++ PATCH] [gomp4] Initial OpenACC support to C++ front-end
- Authentication-results: sourceware.org; auth=none
- References: <5319AED0 dot 3010606 at samsung dot com> <5319AF79 dot 3020808 at samsung dot com>
On 07.03.2014 15:37, Ilmir Usmanov wrote:
Hi Thomas!
I prepared simple patch to add support of OpenACC data, kernels and
parallel constructs to C++ FE.
It adds support of data clauses too.
OK to gomp4 branch?
Fixed subject: changed file extensions of tests and fixed comments.
OK to gomp4 branch?
--
Ilmir.
>From 8368b5196c1201401e1f8301107f11c9e6f064b1 Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usmanov@samsung.com>
Date: Thu, 13 Mar 2014 20:59:34 +0400
Subject: [PATCH] Initial OpenACC support to C++ FE
---
Initial OpenACC support to C++ front-end: parallel, kernels and data
construct with data clauses.
gcc/cp/
* cp-tree.h (finish_oacc_data): New function prototype.
(finish_oacc_kernels, finish_oacc_parallel): Likewise.
* parser.c (cp_parser_omp_clause_name): Support data clauses.
(cp_parser_oacc_data_clause): New function.
(cp_parser_oacc_data_clause_deviceptr, cp_parser_oacc_all_clauses,
cp_parser_oacc_data, cp_parser_oacc_kernels,
cp_parser_oacc_parallel): Likewise.
(OACC_DATA_CLAUSE_MASK): New define.
(OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Likewise.
(cp_parser_omp_construct, cp_parser_pragma): Support OpenACC directives.
* semantics.c (finish_oacc_data): New function.
(finish_oacc_kernels, finish_oacc_parallel): Likewise.
* gcc/testsuite/c-c++-common/goacc/deviceptr-1.c: Move to ...
* gcc/testsuite/gcc.dg/goacc/deviceptr-1.c ... here.
* gcc/testsuite/g++.dg/goacc/goacc.exp: New test directory.
* gcc/testsuite/g++.dg/goacc-gomp/goacc-gomp.exp: Likewise.
gcc/testsuite/g++.dg/goacc/
* deviceptr-1.cpp: New test.
* sb-1.cpp: Likewise.
* sb-2.cpp: Likewise.
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 8ec7d6a..9b966d2 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5825,6 +5825,9 @@ extern tree finish_omp_clauses (tree);
extern void finish_omp_threadprivate (tree);
extern tree begin_omp_structured_block (void);
extern tree finish_omp_structured_block (tree);
+extern tree finish_oacc_data (tree, tree);
+extern tree finish_oacc_kernels (tree, tree);
+extern tree finish_oacc_parallel (tree, tree);
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);
extern tree begin_omp_task (void);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 8c167c7..94d9e22 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -26935,16 +26935,26 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'c':
if (!strcmp ("collapse", p))
result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+ else if (!strcmp ("copy", p))
+ result = PRAGMA_OMP_CLAUSE_COPY;
else if (!strcmp ("copyin", p))
result = PRAGMA_OMP_CLAUSE_COPYIN;
+ else if (!strcmp ("copyout", p))
+ result = PRAGMA_OMP_CLAUSE_COPYOUT;
else if (!strcmp ("copyprivate", p))
result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+ else if (!strcmp ("create", p))
+ result = PRAGMA_OMP_CLAUSE_CREATE;
break;
case 'd':
- if (!strcmp ("depend", p))
+ if (!strcmp ("delete", p))
+ result = PRAGMA_OMP_CLAUSE_DELETE;
+ else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
+ else if (!strcmp ("deviceptr", p))
+ result = PRAGMA_OMP_CLAUSE_DEVICEPTR;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -26993,6 +27003,22 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'p':
if (!strcmp ("parallel", p))
result = PRAGMA_OMP_CLAUSE_PARALLEL;
+ else if (!strcmp ("present", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT;
+ else if (!strcmp ("present_or_copy", p)
+ || !strcmp ("pcopy", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY;
+ else if (!strcmp ("present_or_copyin", p)
+ || !strcmp ("pcopyin", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN;
+ else if (!strcmp ("present_or_copyout", p)
+ || !strcmp ("pcopyout", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT;
+ else if (!strcmp ("present_or_create", p)
+ || !strcmp ("pcreate", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE;
+ else if (!strcmp ("private", p))
+ result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (!strcmp ("proc_bind", p))
result = PRAGMA_OMP_CLAUSE_PROC_BIND;
break;
@@ -27200,6 +27226,111 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
return list;
}
+/* OpenACC 2.0:
+ copy ( variable-list )
+ copyin ( variable-list )
+ copyout ( variable-list )
+ create ( variable-list )
+ delete ( variable-list )
+ present ( variable-list )
+ present_or_copy ( variable-list )
+ pcopy ( variable-list )
+ present_or_copyin ( variable-list )
+ pcopyin ( variable-list )
+ present_or_copyout ( variable-list )
+ pcopyout ( variable-list )
+ present_or_create ( variable-list )
+ pcreate ( variable-list ) */
+
+static tree
+cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
+ tree list)
+{
+ enum omp_clause_map_kind kind;
+ switch (c_kind)
+ {
+ default:
+ gcc_unreachable ();
+ case PRAGMA_OMP_CLAUSE_COPY:
+ kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ kind = OMP_CLAUSE_MAP_FORCE_ALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ kind = OMP_CLAUSE_MAP_FORCE_DEALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ kind = OMP_CLAUSE_MAP_FORCE_PRESENT;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ kind = OMP_CLAUSE_MAP_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ kind = OMP_CLAUSE_MAP_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ kind = OMP_CLAUSE_MAP_ALLOC;
+ break;
+ }
+ tree nl, c;
+ nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
+
+ for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_MAP_KIND (c) = kind;
+
+ return nl;
+}
+
+/* OpenACC 2.0:
+ deviceptr ( variable-list ) */
+
+static tree
+cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
+{
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree vars, t;
+
+ /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic
+ cp_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR,
+ variable-list must only allow for pointer variables. */
+ vars = cp_parser_omp_var_list (parser, OMP_CLAUSE_ERROR, NULL);
+ for (t = vars; t && t; t = TREE_CHAIN (t))
+ {
+ tree v = TREE_PURPOSE (t);
+
+ /* FIXME diagnostics: Ideally we should keep individual
+ locations for all the variables in the var list to make the
+ following errors more precise. Perhaps
+ c_parser_omp_var_list_parens() should construct a list of
+ locations to go along with the var list. */
+
+ if (TREE_CODE (v) != VAR_DECL)
+ error_at (loc, "%qD is not a variable", v);
+ else if (TREE_TYPE (v) == error_mark_node)
+ ;
+ else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+ error_at (loc, "%qD is not a pointer variable", v);
+
+ tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (u) = OMP_CLAUSE_MAP_FORCE_DEVICEPTR;
+ OMP_CLAUSE_DECL (u) = v;
+ OMP_CLAUSE_CHAIN (u) = list;
+ list = u;
+ }
+
+ return list;
+}
+
/* OpenMP 3.0:
collapse ( constant-expression ) */
@@ -28112,6 +28243,101 @@ cp_parser_omp_clause_proc_bind (cp_parser *parser, tree list,
return list;
}
+/* Parse all OpenACC clauses. The set clauses allowed by the directive
+ is a bitmask in MASK. Return the list of clauses found. */
+
+static tree
+cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
+ const char *where, cp_token *pragma_tok,
+ bool finish_p = true)
+{
+ tree clauses = NULL;
+ bool first = true;
+
+ while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+ {
+ location_t here;
+ pragma_omp_clause c_kind;
+ const char *c_name;
+ tree prev = clauses;
+
+ if (!first && cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+
+ here = cp_lexer_peek_token (parser->lexer)->location;
+ c_kind = cp_parser_omp_clause_name (parser);
+
+ switch (c_kind)
+ {
+ case PRAGMA_OMP_CLAUSE_COPY:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "create";
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "delete";
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICEPTR:
+ clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_create";
+ break;
+ default:
+ cp_parser_error (parser, "expected clause");
+ goto saw_error;
+ }
+
+ first = false;
+
+ if (((mask >> c_kind) & 1) == 0)
+ {
+ /* Remove the invalid clause(s) from the list to avoid
+ confusing the rest of the compiler. */
+ clauses = prev;
+ error_at (here, "%qs is not valid for %qs", c_name, where);
+ }
+ }
+
+ saw_error:
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+
+ if (finish_p)
+ return finish_omp_clauses (clauses);
+
+ return clauses;
+}
+
/* Parse all OpenMP clauses. The set clauses allowed by the directive
is a bitmask in MASK. Return the list of clauses found; the result
of clause default goes in *pdefault. */
@@ -28408,6 +28634,113 @@ cp_parser_omp_structured_block (cp_parser *parser)
return finish_omp_structured_block (stmt);
}
+/* OpenACC 2.0:
+ # pragma acc data oacc-data-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (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_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
+
+static tree
+cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_data (clauses, block);
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc kernels oacc-kernels-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_KERNELS_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (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_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
+
+static tree
+cp_parser_oacc_kernels (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
+ "#pragma acc kernels", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_kernels (clauses, block);
+ return stmt;
+}
+
+/* OpenACC 2.0:
+ # pragma acc parallel oacc-parallel-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_PARALLEL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (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_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
+
+static tree
+cp_parser_oacc_parallel (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+ "#pragma acc parallel", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_parallel (clauses, block);
+ return stmt;
+}
/* OpenMP 2.5:
# pragma omp atomic new-line
expression-stmt
@@ -30970,6 +31303,15 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
switch (pragma_tok->pragma_kind)
{
+ case PRAGMA_OACC_DATA:
+ stmt = cp_parser_oacc_data (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_KERNELS:
+ stmt = cp_parser_oacc_kernels (parser, pragma_tok);
+ break;
+ case PRAGMA_OACC_PARALLEL:
+ stmt = cp_parser_oacc_parallel (parser, pragma_tok);
+ break;
case PRAGMA_OMP_ATOMIC:
cp_parser_omp_atomic (parser, pragma_tok);
return;
@@ -31482,6 +31824,9 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_DATA:
+ case PRAGMA_OACC_KERNELS:
+ case PRAGMA_OACC_PARALLEL:
case PRAGMA_OMP_ATOMIC:
case PRAGMA_OMP_CRITICAL:
case PRAGMA_OMP_DISTRIBUTE:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 787eab8..f57007c 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5968,6 +5968,60 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_DATA. */
+
+tree
+finish_oacc_data (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_KERNELS. */
+
+tree
+finish_oacc_kernels (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_KERNELS);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_KERNELS_CLAUSES (stmt) = clauses;
+ OACC_KERNELS_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
+/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_PARALLEL. */
+
+tree
+finish_oacc_parallel (tree clauses, tree block)
+{
+ tree stmt;
+
+ block = finish_omp_structured_block (block);
+
+ stmt = make_node (OACC_PARALLEL);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_PARALLEL_CLAUSES (stmt) = clauses;
+ OACC_PARALLEL_BODY (stmt) = block;
+
+ return add_stmt (stmt);
+}
+
/* Similarly, except force the retention of the BLOCK. */
tree
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
deleted file mode 100644
index 1ac63bd..0000000
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ /dev/null
@@ -1,64 +0,0 @@
-void
-fun1 (void)
-{
-#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */
- ;
-#pragma acc kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
- ;
-
-#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
- ;
-#pragma acc parallel deviceptr(fun1[2:5])
- /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */
- /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */
- ;
-
- int i;
-#pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
- ;
-#pragma acc data deviceptr(i[0:4])
- /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */
- /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */
- ;
-
- float fa[10];
-#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
- ;
-#pragma acc kernels deviceptr(fa[1:5])
- /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */
- /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */
- ;
-
- float *fp;
-#pragma acc data deviceptr(fp)
- ;
-#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
- ;
-}
-
-void
-fun2 (void)
-{
- int i;
- float *fp;
-#pragma acc kernels deviceptr(fp,u,fun2,i,fp)
- /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */
- /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */
- /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */
- /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */
- ;
-}
-
-void
-fun3 (void)
-{
- float *fp;
-#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
- ;
-#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
- ;
-#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
- ;
-}
-
-/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
diff --git a/gcc/testsuite/g++.dg/goacc-gomp/goacc-gomp.exp b/gcc/testsuite/g++.dg/goacc-gomp/goacc-gomp.exp
new file mode 100644
index 0000000..e160901
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc-gomp/goacc-gomp.exp
@@ -0,0 +1,38 @@
+# Copyright (C) 2006-2013 Free Software Foundation, Inc.
+#
+# This file is part of GCC.
+#
+# GCC is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+#
+# GCC is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3. If not see
+# <http://www.gnu.org/licenses/>.
+
+# GCC testsuite that uses the `dg.exp' driver.
+
+# Load support procs.
+load_lib g++-dg.exp
+
+if { ![check_effective_target_fopenacc] \
+ || ![check_effective_target_fopenmp] } {
+ return
+}
+
+# Initialize `dg'.
+dg-init
+
+# Main loop.
+dg-runtest [lsort [concat \
+ [find $srcdir/$subdir *.C] \
+ [find $srcdir/c-c++-common/goacc-gomp *.c]]] "" "-fopenacc -fopenmp"
+
+# All done.
+dg-finish
diff --git a/gcc/testsuite/g++.dg/goacc/deviceptr-1.C b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
new file mode 100644
index 0000000..814af70
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/deviceptr-1.C
@@ -0,0 +1,65 @@
+void
+fun1 (void)
+{
+#pragma acc parallel deviceptr(u) /* { dg-error "'u' has not been declared" } */
+ ;
+#pragma acc kernels deviceptr(u[0:4]) /* { dg-error "'u' has not been declared" } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 6 } */
+ ;
+
+#pragma acc data deviceptr(fun1) /* { dg-error "'void fun1\\\(\\\)' is not a variable" } */
+ ;
+#pragma acc parallel deviceptr(fun1[2:5])
+ /* { dg-error "'void fun1\\\(\\\)' is not a variable" "not a variable" { target *-*-* } 12 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 12 } */
+ ;
+
+ int i;
+#pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
+ ;
+#pragma acc data deviceptr(i[0:4])
+ /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 20 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 20 } */
+ ;
+
+ float fa[10];
+#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
+ ;
+#pragma acc kernels deviceptr(fa[1:5])
+ /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 28 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 28 } */
+ ;
+
+ float *fp;
+#pragma acc data deviceptr(fp)
+ ;
+#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+}
+
+void
+fun2 (void)
+{
+ int i;
+ float *fp;
+#pragma acc kernels deviceptr(fp,u,fun2,i,fp)
+ /* { dg-error "'u' has not been declared" "u undeclared" { target *-*-* } 45 } */
+ /* { dg-error "'void fun2\\\(\\\)' is not a variable" "fun2 not a variable" { target *-*-* } 45 } */
+ /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 45 } */
+ /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 45 } */
+ ;
+}
+
+void
+fun3 (void)
+{
+ float *fp;
+#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+}
+
+/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
diff --git a/gcc/testsuite/g++.dg/goacc/goacc.exp b/gcc/testsuite/g++.dg/goacc/goacc.exp
new file mode 100644
index 0000000..8ccb1fd
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/goacc.exp
@@ -0,0 +1,37 @@
+# Copyright (C) 2006-2013 Free Software Foundation, Inc.
+#
+# This file is part of GCC.
+#
+# GCC is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3, or (at your option)
+# any later version.
+#
+# GCC is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with GCC; see the file COPYING3. If not see
+# <http://www.gnu.org/licenses/>.
+
+# GCC testsuite that uses the `dg.exp' driver.
+
+# Load support procs.
+load_lib g++-dg.exp
+
+if ![check_effective_target_fopenacc] {
+ return
+}
+
+# Initialize `dg'.
+dg-init
+
+# Main loop.
+dg-runtest [lsort [concat \
+ [find $srcdir/$subdir *.C] \
+ [find $srcdir/c-c++-common/goacc *.c]]] "" "-fopenacc"
+
+# All done.
+dg-finish
diff --git a/gcc/testsuite/g++.dg/goacc/sb-1.C b/gcc/testsuite/g++.dg/goacc/sb-1.C
new file mode 100644
index 0000000..6dbb21a
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/sb-1.C
@@ -0,0 +1,59 @@
+// { dg-do compile }
+
+void foo()
+{
+ bad1:
+ #pragma acc parallel
+ goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
+ #pragma acc kernels
+ goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
+ #pragma acc data
+ goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
+ goto bad2_parallel; // { dg-error "invalid entry to OpenACC structured block" }
+// { dg-error "from here" "from here" { target *-*-* } 12 }
+ #pragma acc parallel
+ {
+ bad2_parallel: ; // { dg-error "jump to label" }
+// { dg-error "enters OpenMP structured block" "enters OpenMP structured block" { target *-*-* } 16 }
+ }
+
+ goto bad2_kernels; // { dg-error "invalid entry to OpenACC structured block" }
+// { dg-error "from here" "from here" { target *-*-* } 20 }
+ #pragma acc kernels
+ {
+ bad2_kernels: ; // { dg-error "jump to label" }
+// { dg-error "enters OpenMP structured block" "enters OpenMP structured block" { target *-*-* } 24 }
+ }
+
+ goto bad2_data; // { dg-error "invalid entry to OpenACC structured block" }
+// { dg-error "from here" "from here" { target *-*-* } 28 }
+ #pragma acc data
+ {
+ bad2_data: ; // { dg-error "jump to label" }
+// { dg-error "enters OpenMP structured block" "enters OpenMP structured block" { target *-*-* } 32 }
+ }
+
+ #pragma acc parallel
+ {
+ int i;
+ goto ok1_parallel;
+ for (i = 0; i < 10; ++i)
+ { ok1_parallel: break; }
+ }
+
+ #pragma acc kernels
+ {
+ int i;
+ goto ok1_kernels;
+ for (i = 0; i < 10; ++i)
+ { ok1_kernels: break; }
+ }
+
+ #pragma acc data
+ {
+ int i;
+ goto ok1_data;
+ for (i = 0; i < 10; ++i)
+ { ok1_data: break; }
+ }
+}
diff --git a/gcc/testsuite/g++.dg/goacc/sb-2.C b/gcc/testsuite/g++.dg/goacc/sb-2.C
new file mode 100644
index 0000000..cc36394
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/sb-2.C
@@ -0,0 +1,25 @@
+// { dg-do compile }
+
+void foo(int i)
+{
+ switch (i)
+ {
+ #pragma acc parallel
+ { case 0:; } // { dg-error "enters OpenMP structured block" }
+// { dg-error "jump to case label" "jump to case label" { target *-*-* } 8 }
+ }
+
+ switch (i)
+ {
+ #pragma acc kernels
+ { case 0:; } // { dg-error "enters OpenMP structured block" }
+// { dg-error "jump to case label" "jump to case label" { target *-*-* } 15 }
+ }
+
+ switch (i)
+ {
+ #pragma acc data
+ { case 0:; } // { dg-error "enters OpenMP structured block" }
+// { dg-error "jump to case label" "jump to case label" { target *-*-* } 22 }
+ }
+}
diff --git a/gcc/testsuite/gcc.dg/goacc/deviceptr-1.c b/gcc/testsuite/gcc.dg/goacc/deviceptr-1.c
new file mode 100644
index 0000000..1ac63bd
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/deviceptr-1.c
@@ -0,0 +1,64 @@
+void
+fun1 (void)
+{
+#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */
+ ;
+#pragma acc kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+
+#pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
+ ;
+#pragma acc parallel deviceptr(fun1[2:5])
+ /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */
+ ;
+
+ int i;
+#pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
+ ;
+#pragma acc data deviceptr(i[0:4])
+ /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */
+ ;
+
+ float fa[10];
+#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
+ ;
+#pragma acc kernels deviceptr(fa[1:5])
+ /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */
+ ;
+
+ float *fp;
+#pragma acc data deviceptr(fp)
+ ;
+#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+}
+
+void
+fun2 (void)
+{
+ int i;
+ float *fp;
+#pragma acc kernels deviceptr(fp,u,fun2,i,fp)
+ /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */
+ /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */
+ /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */
+ /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */
+ ;
+}
+
+void
+fun3 (void)
+{
+ float *fp;
+#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+}
+
+/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
--
1.8.3.2