This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[patch,gomp4] error on invalid acc loop clauses
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Cc: Jakub Jelinek <jakub at redhat dot com>, Thomas Schwinge <thomas at codesourcery dot com>
- Date: Fri, 15 May 2015 11:10:21 -0700
- Subject: [patch,gomp4] error on invalid acc loop clauses
- Authentication-results: sourceware.org; auth=none
This patch teaches the c and c++ front ends to error on invalid and
conflicting acc loop clauses. E.g., an acc loop cannot have 'gang seq'
and the worker and vector clauses inside parallel regions cannot have
optional kernel-specific arguments.
The c and c++ front end also error when it detects a parallel or kernels
region nested inside a parallel or kernels region. E.g.
#pragma acc parallel
{
#pragma acc parallel
...
}
This is technically supported by OpenACC 2.0a, but there are a couple of
unresolved technical issues preventing its inclusion for gcc 6.0. For
starters, we don't have the runtime rigged up to handle CUDA's dynamic
parallelism, and it's unclear if the spec itself requires the runtime to
do so. Then there's the issue of mapping gangs within gangs, which
presumably gets handled by dynamic parallelism.
I included two new test cases in this patch. They are mostly identical
but, unfortunately, the c and c++ front ends emit slightly different
error messages.
The front ends still need to be cleaned before this functionality should
be considered for mainline. So for the time being I've applied this
patch to gomp-4_0-branch.
Cesar
2015-05-15 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-parser.c (typedef struct c_parser): Add BOOL_BITFIELDS
oacc_parallel_region and oacc_kernels_region.
(c_parser_oacc_shape_clause): Only use op1 for the static argument
in the gang clause. Check for incompatible clause arguments inside
parallel regions.
(c_parser_oacc_loop): Error on conflicting loop clauses.
(c_parser_oacc_kernels): Error in nested parallel and kernels.
(c_parser_oacc_parallel): Likewise.
gcc/cp/
* parser.h (typedef struct cp_parser): Add bool oacc_parallel_region
and oacc_kernels_region.
* parser.c (cp_parser_oacc_shape_clause): Only use op1 for the static
argument in the gang clause. Check for incompatible clause arguments
inside parallel regions.
(cp_parser_oacc_loop): Error on conflicting loop clauses.
(cp_parser_oacc_parallel_kernels): Error in nested parallel and
kernels.
gcc/
* omp-low.c (scan_omp_for): Remove check for nested parallel and
kernels regions.
gcc/testsuite/
* c-c++-common/goacc/nesting-fail-1.c: Update error messages.
* g++.dg/goacc/loop-4.C: New test.
* gcc.dg/goacc/loop-1.c: New test.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 74adeb8..f508b91 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -234,6 +234,10 @@ typedef struct GTY(()) c_parser {
/* True if we are in a context where the Objective-C "Property attribute"
keywords are valid. */
BOOL_BITFIELD objc_property_attr_context : 1;
+ /* True if we are inside a OpenACC parallel region. */
+ BOOL_BITFIELD oacc_parallel_region : 1;
+ /* True if we are inside a OpenACC kernels region. */
+ BOOL_BITFIELD oacc_kernels_region : 1;
/* Cilk Plus specific parser/lexer information. */
@@ -10839,6 +10843,7 @@ c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind,
mark_exp_read (expr);
require_positive_expr (expr, expr_loc, str);
*op_to_parse = expr;
+ op_to_parse = &op0;
}
while (!c_parser_next_token_is (parser, CPP_CLOSE_PAREN));
c_parser_consume_token (parser);
@@ -10852,6 +10857,17 @@ c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind,
if (op1)
OMP_CLAUSE_OPERAND (c, 1) = op1;
OMP_CLAUSE_CHAIN (c) = list;
+
+ if (parser->oacc_parallel_region && (op0 != NULL || op1 != NULL))
+ {
+ if (c_kind != PRAGMA_OACC_CLAUSE_GANG)
+ c_parser_error (parser, c_kind == PRAGMA_OACC_CLAUSE_WORKER ?
+ "worker clause arguments are not supported in OpenACC parallel regions"
+ : "vector clause arguments are not supported in OpenACC parallel regions");
+ else if (op0 != NULL)
+ c_parser_error (parser, "non-static argument to clause gang");
+ }
+
return c;
}
@@ -12721,7 +12737,10 @@ static tree
c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
omp_clause_mask mask, tree *cclauses)
{
- tree stmt, clauses, block;
+ tree stmt, clauses, block, c;
+ bool gwv = false;
+ bool auto_clause = false;
+ bool seq_clause = false;
strcat (p_name, " loop");
mask |= OACC_LOOP_CLAUSE_MASK;
@@ -12732,6 +12751,33 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
if (cclauses)
clauses = oacc_split_loop_clauses (clauses, cclauses);
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ gwv = true;
+ break;
+ case OMP_CLAUSE_AUTO:
+ auto_clause = true;
+ break;
+ case OMP_CLAUSE_SEQ:
+ seq_clause = true;
+ break;
+ default:
+ ;
+ }
+ }
+
+ if (gwv && auto_clause)
+ c_parser_error (parser, "incompatible use of clause %<auto%>");
+ else if (gwv && seq_clause)
+ c_parser_error (parser, "incompatible use of clause %<seq%>");
+ else if (auto_clause && seq_clause)
+ c_parser_error (parser, "incompatible use of clause %<seq%> and %<auto%>");
+
block = c_begin_compound_stmt (true);
stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
block = c_end_compound_stmt (loc, block, true);
@@ -12774,6 +12820,13 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
strcat (p_name, " kernels");
+ if (parser->oacc_parallel_region || parser->oacc_kernels_region)
+ {
+ c_parser_error (parser, "nested kernels region");
+ }
+
+ parser->oacc_kernels_region = true;
+
mask = OACC_KERNELS_CLAUSE_MASK;
if (c_parser_next_token_is (parser, CPP_NAME))
{
@@ -12787,6 +12840,7 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
block = c_begin_omp_parallel ();
c_parser_oacc_loop (loc, parser, p_name, mask, &kernel_clauses);
stmt = c_finish_oacc_kernels (loc, kernel_clauses, block);
+ parser->oacc_kernels_region = false;
return stmt;
}
}
@@ -12797,6 +12851,7 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
stmt = c_finish_oacc_kernels (loc, clauses, block);
+ parser->oacc_kernels_region = false;
return stmt;
}
@@ -12843,6 +12898,13 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
strcat (p_name, " parallel");
+ if (parser->oacc_parallel_region || parser->oacc_kernels_region)
+ {
+ c_parser_error (parser, "nested parallel region");
+ }
+
+ parser->oacc_parallel_region = true;
+
mask = OACC_PARALLEL_CLAUSE_MASK;
dmask = OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK;
if (c_parser_next_token_is (parser, CPP_NAME))
@@ -12857,6 +12919,7 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
block = c_begin_omp_parallel ();
c_parser_oacc_loop (loc, parser, p_name, mask, ¶llel_clauses);
stmt = c_finish_oacc_parallel (loc, parallel_clauses, block);
+ parser->oacc_parallel_region = false;
return stmt;
}
}
@@ -12866,6 +12929,7 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
stmt = c_finish_oacc_parallel (loc, clauses, block);
+ parser->oacc_parallel_region = false;
return stmt;
}
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index e418ca2..2947bf4 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -28318,6 +28318,7 @@ cp_parser_oacc_shape_clause (cp_parser *parser, pragma_omp_clause c_kind,
mark_exp_read (expr);
require_positive_expr (expr, expr_loc, str);
*op_to_parse = expr;
+ op_to_parse = &op0;
if (cp_lexer_next_token_is (lexer, CPP_COMMA))
cp_lexer_consume_token (lexer);
@@ -28334,6 +28335,17 @@ cp_parser_oacc_shape_clause (cp_parser *parser, pragma_omp_clause c_kind,
if (op1)
OMP_CLAUSE_OPERAND (c, 1) = op1;
OMP_CLAUSE_CHAIN (c) = list;
+
+ if (parser->oacc_parallel_region && (op0 != NULL || op1 != NULL))
+ {
+ if (c_kind != PRAGMA_OACC_CLAUSE_GANG)
+ cp_parser_error (parser, c_kind == PRAGMA_OACC_CLAUSE_WORKER ?
+ "worker clause arguments are not supported in OpenACC parallel regions"
+ : "vector clause arguments are not supported in OpenACC parallel regions");
+ else if (op0 != NULL)
+ cp_parser_error (parser, "non-static argument to clause gang");
+ }
+
return c;
}
@@ -32198,7 +32210,10 @@ static tree
cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
omp_clause_mask mask, tree *cclauses)
{
- tree stmt, clauses, block;
+ tree stmt, clauses, block, c;
+ bool gwv = false;
+ bool auto_clause = false;
+ bool seq_clause = false;
int save;
strcat (p_name, " loop");
@@ -32211,6 +32226,33 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
if (cclauses)
clauses = oacc_split_loop_clauses (clauses, cclauses);
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ gwv = true;
+ break;
+ case OMP_CLAUSE_AUTO:
+ auto_clause = true;
+ break;
+ case OMP_CLAUSE_SEQ:
+ seq_clause = true;
+ break;
+ default:
+ ;
+ }
+ }
+
+ if (gwv && auto_clause)
+ cp_parser_error (parser, "incompatible use of clause %<auto%>");
+ else if (gwv && seq_clause)
+ cp_parser_error (parser, "incompatible use of clause %<seq%>");
+ else if (auto_clause && seq_clause)
+ cp_parser_error (parser, "incompatible use of clause %<seq%> and %<auto%>");
+
block = begin_omp_structured_block ();
save = cp_parser_begin_omp_structured_block (parser);
stmt = cp_parser_omp_for_loop (parser, OACC_LOOP, clauses, NULL);
@@ -32288,13 +32330,21 @@ cp_parser_oacc_parallel_kernels (cp_parser *parser, cp_token *pragma_tok,
cp_lexer *lexer = parser->lexer;
omp_clause_mask mask, dtype_mask;
+ if (parser->oacc_parallel_region || parser->oacc_kernels_region)
+ {
+ cp_parser_error (parser, is_parallel ? "nested parallel region"
+ : "nested kernels region");
+ }
+
if (is_parallel)
{
+ parser->oacc_parallel_region = true;
mask = OACC_PARALLEL_CLAUSE_MASK;
strcat (p_name, " parallel");
}
else
{
+ parser->oacc_kernels_region = true;
mask = OACC_KERNELS_CLAUSE_MASK;
strcat (p_name, " kernels");
}
@@ -32313,6 +32363,10 @@ cp_parser_oacc_parallel_kernels (cp_parser *parser, cp_token *pragma_tok,
&combined_clauses);
stmt = is_parallel ? finish_oacc_parallel (combined_clauses, block)
: finish_oacc_kernels (combined_clauses, block);
+ if (is_parallel)
+ parser->oacc_parallel_region = false;
+ else
+ parser->oacc_kernels_region = false;
return stmt;
}
}
@@ -32329,6 +32383,10 @@ cp_parser_oacc_parallel_kernels (cp_parser *parser, cp_token *pragma_tok,
cp_parser_end_omp_structured_block (parser, save);
stmt = is_parallel ? finish_oacc_parallel (clauses, block)
: finish_oacc_kernels (clauses, block);
+ if (is_parallel)
+ parser->oacc_parallel_region = false;
+ else
+ parser->oacc_kernels_region = false;
return stmt;
}
diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h
index 8eb5484..07292bf 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -377,6 +377,11 @@ typedef struct GTY(()) cp_parser {
cp_omp_declare_simd_data * GTY((skip)) oacc_routine;
vec <tree, va_gc> *named_oacc_routines;
+ /* True if we are inside a OpenACC parallel region. */
+ bool oacc_parallel_region;
+ /* True if we are inside a OpenACC kernels region. */
+ bool oacc_kernels_region;
+
/* Nonzero if parsing a parameter list where 'auto' should trigger an implicit
template parameter. */
bool auto_is_implicit_function_template_parm_p;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index c4f8033..24e8771 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2959,14 +2959,6 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
}
if (outer_type == GIMPLE_OMP_FOR)
outer_ctx->gwv_below |= val;
- if (OMP_CLAUSE_OPERAND (c, 0) != NULL_TREE)
- {
- omp_context *enclosing = enclosing_target_ctx (outer_ctx);
- if (gimple_omp_target_kind (enclosing->stmt)
- == GF_OMP_TARGET_KIND_OACC_PARALLEL)
- error_at (gimple_location (stmt),
- "no arguments allowed to gang, worker and vector clauses inside parallel");
- }
}
}
diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 8af1c82..6054fb8 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -7,9 +7,9 @@ f_acc_parallel (void)
{
#pragma acc parallel
{
-#pragma acc parallel /* { dg-bogus "parallel construct inside of parallel region" "not implemented" { xfail *-*-* } } */
+#pragma acc parallel /* { dg-error "nested parallel region" } */
;
-#pragma acc kernels /* { dg-bogus "kernels construct inside of parallel region" "not implemented" { xfail *-*-* } } */
+#pragma acc kernels /* { dg-error " kernels construct inside of parallel" } */
;
#pragma acc data /* { dg-error "data construct inside of parallel region" } */
;
@@ -26,9 +26,9 @@ f_acc_kernels (void)
{
#pragma acc kernels
{
-#pragma acc parallel /* { dg-bogus "parallel construct inside of kernels region" "not implemented" { xfail *-*-* } } */
+#pragma acc parallel /* { dg-error "nested parallel region" } */
;
-#pragma acc kernels /* { dg-bogus "kernels construct inside of kernels region" "not implemented" { xfail *-*-* } } */
+#pragma acc kernels /* { dg-error "nested kernels region" } */
;
#pragma acc data /* { dg-error "data construct inside of kernels region" } */
;
@@ -37,3 +37,8 @@ f_acc_kernels (void)
#pragma acc exit data delete(i) /* { dg-error "enter/exit data construct inside of kernels region" } */
}
}
+
+// { dg-error "parallel construct inside of parallel" "" { target *-*-* } 10 }
+
+// { dg-error "parallel construct inside of kernels" "" { target *-*-* } 29 }
+// { dg-error "kernels construct inside of kernels" "" { target *-*-* } 31 }
diff --git a/gcc/testsuite/g++.dg/goacc/loop-4.C b/gcc/testsuite/g++.dg/goacc/loop-4.C
new file mode 100644
index 0000000..5361963
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/loop-4.C
@@ -0,0 +1,686 @@
+// { dg-do compile }
+// { dg-options "-fopenacc" }
+// { dg-additional-options "-fmax-errors=200" }
+
+int
+main ()
+{
+ int i, j;
+
+#pragma acc kernels
+ {
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 0; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 0; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected primary-expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(6-2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(6+2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 0; j < 10; i++)
+ { }
+ }
+#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2, 2, 1)
+ for (i = 2; i < 4; i++)
+ for (i = 4; i < 6; i++)
+ { }
+#pragma acc loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ for (j = i+1; j < 7; i++)
+ { }
+#pragma acc loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+ }
+
+
+#pragma acc parallel
+ {
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(5) // { dg-error "non-static" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5) // { dg-error "non-static" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5) // { dg-error "worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5) // { dg-error "worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected primary-expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ {
+ for (j = i + 1; j < 7; j+=i)
+ { }
+ }
+#pragma acc loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+ }
+
+#pragma acc kernels loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop gang // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc kernels loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop worker // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc kernels loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector(length:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop vector // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc kernels loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc kernels loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile() // { dg-error "expected primary-expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc kernels loop tile(2, 2)
+ for (i = 1; i < 5; i++)
+ {
+ for (j = i + 1; j < 7; j += i)
+ { }
+ }
+#pragma acc kernels loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(5) // { dg-error "non-static argument to clause gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(num:5) // { dg-error "non-static argument to clause gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop gang // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(5) // { dg-error " worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(num:5) // { dg-error " worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop worker // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector(5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector(length:5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop vector // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc parallel loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile() // { dg-error "expected primary-expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc parallel loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ {
+ for (j = i + 1; j < 7; j++)
+ { }
+ }
+#pragma acc parallel loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+
+ return 0;
+}
+
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 377 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 397 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 400 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 423 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 426 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 429 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 535 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 555 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 558 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 581 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 584 }
+// { dg-error "OpenACC construct inside of non-OpenACC region" "" { target *-*-* } 587 }
diff --git a/gcc/testsuite/gcc.dg/goacc/loop-1.c b/gcc/testsuite/gcc.dg/goacc/loop-1.c
new file mode 100644
index 0000000..29dcdc5
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/loop-1.c
@@ -0,0 +1,674 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fmax-errors=200" } */
+
+int
+main ()
+{
+ int i, j;
+
+#pragma acc kernels
+ {
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 0; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 0; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 0; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(6-2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(6+2)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 0; j < 10; i++)
+ { }
+ }
+#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2, 2, 1)
+ for (i = 2; i < 4; i++)
+ for (i = 4; i < 6; i++)
+ { }
+#pragma acc loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ for (j = i+1; j < 7; i++)
+ { }
+#pragma acc loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+ }
+
+
+#pragma acc parallel
+ {
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(5) // { dg-error "non-static" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5) // { dg-error "non-static" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5) // { dg-error "worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5) // { dg-error "worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc loop vector
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile(2)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ {
+ for (j = i + 1; j < 7; j+=i)
+ { }
+ }
+#pragma acc loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+ }
+
+#pragma acc kernels loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop gang // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc kernels loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker(num:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop worker // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc kernels loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector(5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector(length:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc kernels loop vector // { dg-error "nested kernels region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc kernels loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc kernels loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc kernels loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile() // { dg-error "expected expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc kernels loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc kernels loop tile(2, 2)
+ for (i = 1; i < 5; i++)
+ {
+ for (j = i + 1; j < 7; j += i)
+ { }
+ }
+#pragma acc kernels loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(5) // { dg-error "non-static argument to clause gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(num:5) // { dg-error "non-static argument to clause gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(static:5)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang(static:*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop gang // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq gang
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(5) // { dg-error " worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(num:5) // { dg-error " worker clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop worker // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq worker
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang worker
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector(5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector(length:5) // { dg-error "vector clause arguments are not supported" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop vector // { dg-error "nested parallel region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop worker
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq vector
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang vector
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker vector
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc parallel loop auto
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop seq auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop gang auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop worker auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+#pragma acc parallel loop vector auto
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" }
+ { }
+
+#pragma acc parallel loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile() // { dg-error "expected expression" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(1)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(*, 1)
+ for (i = 0; i < 10; i++)
+ {
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+ for (i = 1; i < 10; i++)
+ { }
+#pragma acc parallel loop tile(2, 2, 1)
+ for (i = 1; i < 3; i++)
+ {
+ for (j = 4; j < 6; j++)
+ { }
+ }
+#pragma acc parallel loop tile(2, 2)
+ for (i = 1; i < 5; i+=2)
+ {
+ for (j = i + 1; j < 7; j++)
+ { }
+ }
+#pragma acc parallel loop vector tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector gang tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop vector worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop gang worker tile(*)
+ for (i = 0; i < 10; i++)
+ { }
+
+ return 0;
+}
+
+// { dg-excess-errors "invalid controlling predicate" }