@@ -1,3 +1,8 @@
+2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add "kernels".
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_KERNELS.
+
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
* c-pragma.c (oacc_pragmas): Add "data".
@@ -1170,6 +1170,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 },
+ { "kernels", PRAGMA_OACC_KERNELS },
{ "parallel", PRAGMA_OACC_PARALLEL },
};
static const struct omp_pragma_def omp_pragmas[] = {
@@ -28,6 +28,7 @@ typedef enum pragma_kind {
PRAGMA_NONE = 0,
PRAGMA_OACC_DATA,
+ PRAGMA_OACC_KERNELS,
PRAGMA_OACC_PARALLEL,
PRAGMA_OMP_ATOMIC,
PRAGMA_OMP_BARRIER,
@@ -1,3 +1,11 @@
+2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (OACC_KERNELS_CLAUSE_MASK): New macro definition.
+ (c_parser_oacc_kernels): New function.
+ (c_parser_omp_construct): Handle PRAGMA_OACC_KERNELS.
+ * c-tree.h (c_finish_oacc_kernels): New prototype.
+ * c-typeck.c (c_finish_oacc_kernels): New function.
+
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition.
@@ -4776,11 +4776,15 @@ c_parser_label (c_parser *parser)
openacc-construct:
parallel-construct
+ kernels-construct
data-construct
parallel-construct:
parallel-directive structured-block
+ kernels-construct:
+ kernels-directive structured-block
+
data-construct:
data-directive structured-block
@@ -11401,6 +11405,41 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
}
/* 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
+c_parser_oacc_kernels (location_t loc, c_parser *parser)
+{
+ tree stmt, clauses, block;
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
+ "#pragma acc kernels");
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_kernels (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
# pragma acc parallel oacc-parallel-clause[optseq] new-line
structured-block
@@ -13717,6 +13756,9 @@ c_parser_omp_construct (c_parser *parser)
case PRAGMA_OACC_DATA:
stmt = c_parser_oacc_data (loc, parser);
break;
+ case PRAGMA_OACC_KERNELS:
+ stmt = c_parser_oacc_kernels (loc, parser);
+ break;
case PRAGMA_OACC_PARALLEL:
stmt = c_parser_oacc_parallel (loc, parser);
break;
@@ -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_kernels (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);
@@ -11122,6 +11122,25 @@ c_finish_oacc_parallel (location_t loc, tree clauses, tree 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
+c_finish_oacc_kernels (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_KERNELS);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_KERNELS_CLAUSES (stmt) = clauses;
+ OACC_KERNELS_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
statement. LOC is the location of the OACC_DATA. */
@@ -1,3 +1,19 @@
+2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
+ kernels construct.
+ * c-c++-common/goacc/clauses-fail.c: Likewise.
+ * c-c++-common/goacc/data-clause-duplicate-1.c: Likewise.
+ * c-c++-common/goacc/deviceptr-1.c: Likewise.
+ * c-c++-common/goacc/nesting-fail-1.c: Likewise.
+ * c-c++-common/goacc/kernels-1.c: New file.
+ * gcc.dg/goacc/parallel-sb-1.c: Rename to...
+ * gcc.dg/goacc/sb-1.c: ... this new file, and extend for OpenACC
+ kernels and data constructs.
+ * gcc.dg/goacc/parallel-sb-2.c: Rename to...
+ * gcc.dg/goacc/sb-2.c: ... this new file, and extend for OpenACC
+ kernels and data constructs.
+
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
@@ -9,6 +9,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -18,6 +20,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -30,6 +34,11 @@ f_omp (void)
}
#pragma omp section
{
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
+ }
+#pragma omp section
+ {
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -39,6 +48,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -47,6 +58,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -55,6 +68,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -63,6 +78,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -71,6 +88,8 @@ f_omp (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -144,6 +163,71 @@ f_acc_parallel (void)
/* TODO: Some of these should either be allowed or fail with a more sensible
error message. */
void
+f_acc_kernels (void)
+{
+#pragma acc kernels
+ {
+#pragma omp parallel /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc kernels
+ {
+ int i;
+#pragma omp for /* { dg-error "may not be nested" } */
+ for (i = 0; i < 3; i++)
+ ;
+ }
+
+#pragma acc kernels
+ {
+#pragma omp sections /* { dg-error "may not be nested" } */
+ {
+ ;
+ }
+ }
+
+#pragma acc kernels
+ {
+#pragma omp single /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc kernels
+ {
+#pragma omp task /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc kernels
+ {
+#pragma omp master /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc kernels
+ {
+#pragma omp critical /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc kernels
+ {
+ int i;
+#pragma omp atomic write
+ i = 0; /* { dg-error "may not be nested" } */
+ }
+
+#pragma acc kernels
+ {
+#pragma omp ordered /* { dg-error "may not be nested" } */
+ ;
+ }
+}
+
+/* TODO: Some of these should either be allowed or fail with a more sensible
+ error message. */
+void
f_acc_data (void)
{
#pragma acc data
@@ -4,6 +4,9 @@ f (void)
#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
;
+#pragma acc kernels eins /* { dg-error "expected clause before 'eins'" } */
+ ;
+
#pragma acc data two /* { dg-error "expected clause before 'two'" } */
;
}
@@ -4,9 +4,9 @@ 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 kernels 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)
+#pragma acc data 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 } */
;
@@ -3,10 +3,10 @@ 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 kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
;
-#pragma acc parallel deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
+#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 } */
@@ -14,9 +14,9 @@ fun1 (void)
;
int i;
-#pragma acc parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
+#pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
;
-#pragma acc parallel deviceptr(i[0:4])
+#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 } */
;
@@ -24,13 +24,13 @@ fun1 (void)
float fa[10];
#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
;
-#pragma acc parallel deviceptr(fa[1:5])
+#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 parallel deviceptr(fp)
+#pragma acc data deviceptr(fp)
;
#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
;
@@ -41,7 +41,7 @@ fun2 (void)
{
int i;
float *fp;
-#pragma acc parallel deviceptr(fp,u,fun2,i,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 } */
@@ -53,11 +53,11 @@ void
fun3 (void)
{
float *fp;
-#pragma acc parallel deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#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 parallel copy(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" } */
;
}
new file mode 100644
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc kernels
+ ;
+}
@@ -7,6 +7,24 @@ f_acc_parallel (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { 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_kernels (void)
+{
+#pragma acc parallel
+ {
+#pragma acc parallel /* { dg-error "may not be nested" } */
+ ;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
@@ -21,6 +39,8 @@ f_acc_data (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc kernels /* { dg-error "may not be nested" } */
+ ;
#pragma acc data /* { dg-error "may not be nested" } */
;
}
deleted file mode 100644
@@ -1,22 +0,0 @@
-// { dg-do compile }
-
-void foo()
-{
- bad1:
- #pragma acc parallel
- goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
-
- goto bad2; // { dg-error "invalid entry to OpenACC structured block" }
- #pragma acc parallel
- {
- bad2: ;
- }
-
- #pragma acc parallel
- {
- int i;
- goto ok1;
- for (i = 0; i < 10; ++i)
- { ok1: break; }
- }
-}
deleted file mode 100644
@@ -1,10 +0,0 @@
-// { dg-do compile }
-
-void foo(int i)
-{
- switch (i) // { dg-error "invalid entry to OpenACC structured block" }
- {
- #pragma acc parallel
- { case 0:; }
- }
-}
new file mode 100644
@@ -0,0 +1,54 @@
+// { 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" }
+ #pragma acc parallel
+ {
+ bad2_parallel: ;
+ }
+
+ goto bad2_kernels; // { dg-error "invalid entry to OpenACC structured block" }
+ #pragma acc kernels
+ {
+ bad2_kernels: ;
+ }
+
+ goto bad2_data; // { dg-error "invalid entry to OpenACC structured block" }
+ #pragma acc data
+ {
+ bad2_data: ;
+ }
+
+ #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; }
+ }
+}
new file mode 100644
@@ -0,0 +1,22 @@
+// { dg-do compile }
+
+void foo(int i)
+{
+ switch (i) // { dg-error "invalid entry to OpenACC structured block" }
+ {
+ #pragma acc parallel
+ { case 0:; }
+ }
+
+ switch (i) // { dg-error "invalid entry to OpenACC structured block" }
+ {
+ #pragma acc kernels
+ { case 0:; }
+ }
+
+ switch (i) // { dg-error "invalid entry to OpenACC structured block" }
+ {
+ #pragma acc data
+ { case 0:; }
+ }
+}
@@ -1,5 +1,9 @@
2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c/goacc_kernels.c: New file.
+ * testsuite/libgomp.oacc-c/kernels-1.c: Likewise.
+ * testsuite/libgomp.oacc-c/parallel-1.c: Add one missing test.
+
* libgomp.map (GOACC_2.0): Add GOACC_kernels.
* libgomp_g.h (GOACC_kernels): New prototype.
* oacc-parallel.c (GOACC_kernels): New function.
new file mode 100644
@@ -0,0 +1,25 @@
+/* { dg-do run } */
+
+#include "libgomp_g.h"
+
+extern void abort ();
+
+volatile int i;
+
+void
+f (void *data)
+{
+ if (i != -1)
+ abort ();
+ i = 42;
+}
+
+int main(void)
+{
+ i = -1;
+ GOACC_kernels (0, f, (const void *) 0, 0, (void *) 0, (void *) 0, (void *) 0);
+ if (i != 42)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,170 @@
+/* { dg-do run } */
+
+extern void abort ();
+
+int i;
+
+int main(void)
+{
+ int j, v;
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc kernels /* copyout */ present_or_copyout (v) 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc kernels /* copyout */ present_or_copyout (v) present_or_create (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc kernels /* 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 kernels /* 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;
+}
@@ -116,6 +116,20 @@ int main(void)
if (v != 1 || i != 2 || j != 1)
abort ();
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_create (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+
#if 0
i = -1;
j = -2;
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/c-family/ * c-pragma.c (oacc_pragmas): Add "kernels". * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_KERNELS. gcc/c/ * c-parser.c (OACC_KERNELS_CLAUSE_MASK): New macro definition. (c_parser_oacc_kernels): New function. (c_parser_omp_construct): Handle PRAGMA_OACC_KERNELS. * c-tree.h (c_finish_oacc_kernels): New prototype. * c-typeck.c (c_finish_oacc_kernels): New function. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC kernels construct. * c-c++-common/goacc/clauses-fail.c: Likewise. * c-c++-common/goacc/data-clause-duplicate-1.c: Likewise. * c-c++-common/goacc/deviceptr-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Likewise. * c-c++-common/goacc/kernels-1.c: New file. * gcc.dg/goacc/parallel-sb-1.c: Rename to... * gcc.dg/goacc/sb-1.c: ... this new file, and extend for OpenACC kernels and data constructs. * gcc.dg/goacc/parallel-sb-2.c: Rename to... * gcc.dg/goacc/sb-2.c: ... this new file, and extend for OpenACC kernels and data constructs. libgomp/ * testsuite/libgomp.oacc-c/goacc_kernels.c: New file. * testsuite/libgomp.oacc-c/kernels-1.c: Likewise. * testsuite/libgomp.oacc-c/parallel-1.c: Add one missing test. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208216 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 | 16 ++ .../c-c++-common/goacc-gomp/nesting-fail-1.c | 84 ++++++++++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 3 + .../c-c++-common/goacc/data-clause-duplicate-1.c | 4 +- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 18 +-- gcc/testsuite/c-c++-common/goacc/kernels-1.c | 6 + gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 20 +++ gcc/testsuite/gcc.dg/goacc/parallel-sb-1.c | 22 --- gcc/testsuite/gcc.dg/goacc/parallel-sb-2.c | 10 -- gcc/testsuite/gcc.dg/goacc/sb-1.c | 54 +++++++ gcc/testsuite/gcc.dg/goacc/sb-2.c | 22 +++ libgomp/ChangeLog.gomp | 4 + libgomp/testsuite/libgomp.oacc-c/goacc_kernels.c | 25 +++ libgomp/testsuite/libgomp.oacc-c/kernels-1.c | 170 +++++++++++++++++++++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 14 ++ 22 files changed, 506 insertions(+), 43 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-1.c delete mode 100644 gcc/testsuite/gcc.dg/goacc/parallel-sb-1.c delete mode 100644 gcc/testsuite/gcc.dg/goacc/parallel-sb-2.c create mode 100644 gcc/testsuite/gcc.dg/goacc/sb-1.c create mode 100644 gcc/testsuite/gcc.dg/goacc/sb-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c/goacc_kernels.c create mode 100644 libgomp/testsuite/libgomp.oacc-c/kernels-1.c