From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/c-family/
* c-pragma.c (oacc_pragmas): Add "data".
* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA.
gcc/c/
* c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition.
(c_parser_oacc_data): New function.
(c_parser_omp_construct): Handle PRAGMA_OACC_DATA.
* c-tree.h (c_finish_oacc_data): New prototype.
* c-typeck.c (c_finish_oacc_data): New function.
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
data construct.
* c-c++-common/goacc/nesting-fail-1.c: Likewise.
* c-c++-common/goacc/parallel-fail-1.c: Rename to...
* c-c++-common/goacc/clauses-fail.c: ... this new file. Extend
for OpenACC data construct.
* c-c++-common/goacc/data-1.c: New file.
libgomp/
* testsuite/libgomp.oacc-c/data-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208017
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/c-family/ChangeLog.gomp | 5 +
gcc/c-family/c-pragma.c | 1 +
gcc/c-family/c-pragma.h | 1 +
gcc/c/ChangeLog.gomp | 8 +
gcc/c/c-parser.c | 42 +++++
gcc/c/c-tree.h | 1 +
gcc/c/c-typeck.c | 19 +++
gcc/testsuite/ChangeLog.gomp | 10 ++
.../c-c++-common/goacc-gomp/nesting-fail-1.c | 92 ++++++++++-
gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 9 ++
gcc/testsuite/c-c++-common/goacc/data-1.c | 6 +
gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 18 ++-
gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c | 6 -
libgomp/ChangeLog.gomp | 2 +
libgomp/testsuite/libgomp.oacc-c/data-1.c | 170 +++++++++++++++++++++
15 files changed, 380 insertions(+), 10 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/clauses-fail.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/data-1.c
delete mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c/data-1.c
diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index e092d53..3da377f 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-02-21 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add "data".
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA.
+
2014-01-28 Thomas Schwinge <tho...@codesourcery.com>
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY,
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index f69486a..08374aa 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1169,6 +1169,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
struct omp_pragma_def { const char *name; unsigned int id; };
static const struct omp_pragma_def oacc_pragmas[] = {
+ { "data", PRAGMA_OACC_DATA },
{ "parallel", PRAGMA_OACC_PARALLEL },
};
static const struct omp_pragma_def omp_pragmas[] = {
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index 1ea5b1d..d092f9f 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see
typedef enum pragma_kind {
PRAGMA_NONE = 0,
+ PRAGMA_OACC_DATA,
PRAGMA_OACC_PARALLEL,
PRAGMA_OMP_ATOMIC,
PRAGMA_OMP_BARRIER,
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index b199957..9b95725 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2014-02-21 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition.
+ (c_parser_oacc_data): New function.
+ (c_parser_omp_construct): Handle PRAGMA_OACC_DATA.
+ * c-tree.h (c_finish_oacc_data): New prototype.
+ * c-typeck.c (c_finish_oacc_data): New function.
+
2014-02-17 Thomas Schwinge <tho...@codesourcery.com>
* c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 7850eab..4643722 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -4776,10 +4776,14 @@ c_parser_label (c_parser *parser)
openacc-construct:
parallel-construct
+ data-construct
parallel-construct:
parallel-directive structured-block
+ data-construct:
+ data-directive structured-block
+
OpenMP:
statement:
@@ -11362,6 +11366,41 @@ c_parser_omp_structured_block (c_parser *parser)
}
/* 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
+c_parser_oacc_data (location_t loc, c_parser *parser)
+{
+ tree stmt, clauses, block;
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data");
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_data (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
# pragma acc parallel oacc-parallel-clause[optseq] new-line
structured-block
@@ -13675,6 +13714,9 @@ c_parser_omp_construct (c_parser *parser)
switch (p_kind)
{
+ case PRAGMA_OACC_DATA:
+ stmt = c_parser_oacc_data (loc, parser);
+ break;
case PRAGMA_OACC_PARALLEL:
stmt = c_parser_oacc_parallel (loc, parser);
break;
diff --git gcc/c/c-tree.h gcc/c/c-tree.h
index c174c7a..c84d3d7 100644
--- gcc/c/c-tree.h
+++ gcc/c/c-tree.h
@@ -634,6 +634,7 @@ extern tree c_finish_goto_label (location_t, tree);
extern tree c_finish_goto_ptr (location_t, tree);
extern tree c_expr_to_decl (tree, bool *, bool *);
extern tree c_finish_oacc_parallel (location_t, tree, tree);
+extern tree c_finish_oacc_data (location_t, tree, tree);
extern tree c_begin_omp_parallel (void);
extern tree c_finish_omp_parallel (location_t, tree, tree);
extern tree c_begin_omp_task (void);
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 76d655b..8c4445b 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11122,6 +11122,25 @@ c_finish_oacc_parallel (location_t loc, tree clauses,
tree block)
return add_stmt (stmt);
}
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_DATA. */
+
+tree
+c_finish_oacc_data (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
/* Like c_begin_compound_stmt, except force the retention of the BLOCK. */
tree
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index fbccfa3..41d73b6 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,13 @@
+2014-02-21 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
+ data construct.
+ * c-c++-common/goacc/nesting-fail-1.c: Likewise.
+ * c-c++-common/goacc/parallel-fail-1.c: Rename to...
+ * c-c++-common/goacc/clauses-fail.c: ... this new file. Extend
+ for OpenACC data construct.
+ * c-c++-common/goacc/data-1.c: New file.
+
2014-02-18 Thomas Schwinge <tho...@codesourcery.com>
* gcc.dg/goacc/parallel-sb-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 875ec66..78fb45b 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -1,7 +1,7 @@
/* TODO: Some of these should either be allowed or fail with a more sensible
error message. */
void
-f1 (void)
+f_omp (void)
{
int i;
@@ -9,6 +9,8 @@ f1 (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp for
@@ -16,49 +18,68 @@ f1 (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp sections
{
+ {
#pragma acc parallel /* { dg-error "may not be nested" } */
- ;
+ ;
+ }
+#pragma omp section
+ {
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
+ }
}
#pragma omp single
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp task
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp master
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp critical
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp ordered
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
}
/* TODO: Some of these should either be allowed or fail with a more sensible
error message. */
void
-f2 (void)
+f_acc_parallel (void)
{
#pragma acc parallel
{
@@ -119,3 +140,68 @@ f2 (void)
;
}
}
+
+/* TODO: Some of these should either be allowed or fail with a more sensible
+ error message. */
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma omp parallel /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+ int i;
+#pragma omp for /* { dg-error "may not be nested" } */
+ for (i = 0; i < 3; i++)
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp sections /* { dg-error "may not be nested" } */
+ {
+ ;
+ }
+ }
+
+#pragma acc data
+ {
+#pragma omp single /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp task /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp master /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp critical /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+ int i;
+#pragma omp atomic write
+ i = 0; /* { dg-error "may not be nested" } */
+ }
+
+#pragma acc data
+ {
+#pragma omp ordered /* { dg-error "may not be nested" } */
+ ;
+ }
+}
diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c
gcc/testsuite/c-c++-common/goacc/clauses-fail.c
new file mode 100644
index 0000000..b0dd042
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -0,0 +1,9 @@
+void
+f (void)
+{
+#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
+ ;
+
+#pragma acc data two /* { dg-error "expected clause before 'two'" } */
+ ;
+}
diff --git gcc/testsuite/c-c++-common/goacc/data-1.c
gcc/testsuite/c-c++-common/goacc/data-1.c
new file mode 100644
index 0000000..8094575
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/data-1.c
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc data
+ ;
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 6501397..24a4c11 100644
--- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -1,11 +1,27 @@
/* TODO: While the OpenACC specification does allow for certain kinds of
nesting, we don't support that yet. */
void
-f1 (void)
+f_acc_parallel (void)
{
#pragma acc parallel
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
+ }
+}
+
+/* TODO: While the OpenACC specification does allow for certain kinds of
+ nesting, we don't support that yet. */
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma acc parallel /* { dg-error "may not be nested" } */
+ ;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
}
diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
deleted file mode 100644
index efc6f14..0000000
--- gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
+++ /dev/null
@@ -1,6 +0,0 @@
-void
-foo (void)
-{
-#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" }
*/
- foo ();
-}
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 5c15656..b90b09b 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,7 @@
2014-02-21 Thomas Schwinge <tho...@codesourcery.com>
+ * testsuite/libgomp.oacc-c/data-1.c: New file.
+
* libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
* libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
* oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
diff --git libgomp/testsuite/libgomp.oacc-c/data-1.c
libgomp/testsuite/libgomp.oacc-c/data-1.c
new file mode 100644
index 0000000..8f9a17a
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/data-1.c
@@ -0,0 +1,170 @@
+/* { dg-do run } */
+
+extern void abort ();
+
+int i;
+
+int main(void)
+{
+ int j;
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data copyin (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data copyout (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data copy (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data create (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copyin (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != 2 || j != 1)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copyout (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copy (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present_or_create (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ return 0;
+}
--
1.8.1.1
