[gomp4 6/6] Enable initial support in the C front end for OpenACC data clauses.
thomas@codesourcery.com
thomas@codesourcery.com
Tue Jan 14 15:10:00 GMT 2014
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
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
More information about the Gcc-patches
mailing list