@@ -11225,8 +11225,17 @@ c_parser_omp_structured_block (c_parser *parser)
LOC is the location of the #pragma token.
*/
-#define OACC_PARALLEL_CLAUSE_MASK \
- (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE)
+#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
c_parser_oacc_parallel (location_t loc, c_parser *parser)
@@ -11235,7 +11244,6 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser)
clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
"#pragma acc parallel");
- gcc_assert (clauses == NULL);
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
new file mode 100644
@@ -0,0 +1,13 @@
+void
+fun (void)
+{
+ float *fp;
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel create(fp[:10]) deviceptr(fp)
+ /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */
+ /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */
+ ;
+}
new file mode 100644
@@ -0,0 +1,64 @@
+void
+fun1 (void)
+{
+#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */
+ ;
+#pragma acc parallel deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+
+#pragma acc parallel 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 parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
+ ;
+#pragma acc parallel 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 parallel 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 parallel deviceptr(fp)
+ ;
+#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+}
+
+void
+fun2 (void)
+{
+ int i;
+ float *fp;
+#pragma acc parallel 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 parallel 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 parallel copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+}
+
+/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
@@ -2,25 +2,155 @@
extern void abort ();
-volatile int i;
+int i;
int main(void)
{
- volatile int j;
+ int j, v;
- i = -0x42;
- j = -42;
-#pragma acc parallel
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copyin (i, j)
{
- if (i != -0x42 || j != -42)
+ if (i != -1 || j != -2)
abort ();
- i = 42;
- j = 0x42;
- if (i != 42 || j != 0x42)
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
abort ();
+ v = 1;
}
- if (i != 42 || j != 0x42)
+ if (v != 1 || i != -1 || j != -2)
abort ();
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copyout (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copy (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) create (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyin (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyout (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copy (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+#endif
+
return 0;
}
From: Thomas Schwinge <thomas@codesourcery.com> gcc/c/ * c-parser.c (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: New file. * c-c++-common/goacc/deviceptr-1.c: New file. libgomp/ * testsuite/libgomp.oacc-c/parallel-1.c: Extend. --- gcc/c/c-parser.c | 14 +- .../c-c++-common/goacc/data-clause-duplicate-1.c | 13 ++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 64 +++++++++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 150 +++++++++++++++++++-- 4 files changed, 228 insertions(+), 13 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-1.c