From: Thomas Schwinge <tho...@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
diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 48c55e6..d6a2af0 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -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)); diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c new file mode 100644 index 0000000..1bcf5be --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -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 } */ + ; +} diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c new file mode 100644 index 0000000..0f0cf0c --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -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" } */ diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c index b40545d..ff54b9d 100644 --- libgomp/testsuite/libgomp.oacc-c/parallel-1.c +++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c @@ -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; } -- 1.8.1.1