* [patch,gomp4] error on invalid acc loop clauses
@ 2015-05-15 18:14 ` Cesar Philippidis
2015-05-20 7:48 ` [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used Thomas Schwinge
2015-05-20 8:24 ` [patch,gomp4] error on invalid acc loop clauses Thomas Schwinge
0 siblings, 2 replies; 11+ messages in thread
From: Cesar Philippidis @ 2015-05-15 18:14 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 1176 bytes --]
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
[-- Attachment #2: gwv-loops.diff --]
[-- Type: text/x-patch, Size: 51230 bytes --]
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" }
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used
[not found] <555A05BA.3050001@mentor.com>
@ 2015-05-19 8:04 ` Tom de Vries
2015-05-15 18:14 ` [patch,gomp4] error on invalid acc loop clauses Cesar Philippidis
0 siblings, 1 reply; 11+ messages in thread
From: Tom de Vries @ 2015-05-19 8:04 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: GCC Patches
[ moved to gcc-patches ml ]
On 18-05-15 17:31, Tom de Vries wrote:
> Thomas,
>
> In ran into this bootstrap failure with branch gomp-4_0-branch:
> ...
> src/gcc-gomp-4_0-branch/gcc/omp-low.c:2897:1: error: 'omp_context*
> enclosing_target_ctx(omp_context*)' defined but not used [-Werror=unused-function]
> enclosing_target_ctx (omp_context *ctx)
> ^
> cc1plus: all warnings being treated as errors
> make[3]: *** [omp-low.o] Error 1
> ...
>
This patch fixes bootstrap by commenting out the unused function
enclosing_target_ctx.
The patch just comments it out, since I'm not sure whether:
- the function needs to be removed, or
- a user of the function will soon be committed.
Committed to fix bootstrap.
Thanks,
- Tom
Comment out unused enclosing_target_ctx
2015-05-19 Tom de Vries <tom@codesourcery.com>
* omp-low.c (enclosing_target_ctx): Comment out.
---
gcc/omp-low.c | 2 ++
1 file changed, 2 insertions(+)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 914549c..3414ab5 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2893,6 +2893,7 @@ finish_taskreg_scan (omp_context *ctx)
}
+#if 0
static omp_context *
enclosing_target_ctx (omp_context *ctx)
{
@@ -2902,6 +2903,7 @@ enclosing_target_ctx (omp_context *ctx)
gcc_assert (ctx != NULL);
return ctx;
}
+#endif
static bool
oacc_loop_or_target_p (gimple stmt)
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used
2015-05-15 18:14 ` [patch,gomp4] error on invalid acc loop clauses Cesar Philippidis
@ 2015-05-20 7:48 ` Thomas Schwinge
2015-05-20 8:24 ` [patch,gomp4] error on invalid acc loop clauses Thomas Schwinge
1 sibling, 0 replies; 11+ messages in thread
From: Thomas Schwinge @ 2015-05-20 7:48 UTC (permalink / raw)
To: Tom de Vries, Cesar Philippidis; +Cc: GCC Patches
[-- Attachment #1: Type: text/plain, Size: 1961 bytes --]
Hi!
On Tue, 19 May 2015 09:24:51 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> On 18-05-15 17:31, Tom de Vries wrote:
> > In ran into this bootstrap failure with branch gomp-4_0-branch:
> > ...
> > src/gcc-gomp-4_0-branch/gcc/omp-low.c:2897:1: error: 'omp_context*
> > enclosing_target_ctx(omp_context*)' defined but not used [-Werror=unused-function]
> > enclosing_target_ctx (omp_context *ctx)
> > ^
> > cc1plus: all warnings being treated as errors
> > make[3]: *** [omp-low.o] Error 1
> > ...
I can only encourage everyone to pay attention to compiler warnings.
> This patch fixes bootstrap by commenting out the unused function
> enclosing_target_ctx.
>
> The patch just comments it out, since I'm not sure whether:
> - the function needs to be removed, or
> - a user of the function will soon be committed.
Well, looking at the recent revision history, I see that in r223222 Cesar
has removed the single use of enclosing_target_ctx,
<http://news.gmane.org/find-root.php?message_id=%3C5556368D.7010904%40codesourcery.com%3E>,
so I'd assume it is no longer needed? That is, Cesar, please remove the
function in this case.
> Committed to fix bootstrap.
Thanks!
> Comment out unused enclosing_target_ctx
>
> 2015-05-19 Tom de Vries <tom@codesourcery.com>
>
> * omp-low.c (enclosing_target_ctx): Comment out.
> ---
> gcc/omp-low.c | 2 ++
> 1 file changed, 2 insertions(+)
>
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index 914549c..3414ab5 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -2893,6 +2893,7 @@ finish_taskreg_scan (omp_context *ctx)
> }
>
>
> +#if 0
> static omp_context *
> enclosing_target_ctx (omp_context *ctx)
> {
> @@ -2902,6 +2903,7 @@ enclosing_target_ctx (omp_context *ctx)
> gcc_assert (ctx != NULL);
> return ctx;
> }
> +#endif
>
> static bool
> oacc_loop_or_target_p (gimple stmt)
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-15 18:14 ` [patch,gomp4] error on invalid acc loop clauses Cesar Philippidis
2015-05-20 7:48 ` [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used Thomas Schwinge
@ 2015-05-20 8:24 ` Thomas Schwinge
2015-05-20 8:56 ` Jakub Jelinek
2015-05-20 14:19 ` Cesar Philippidis
1 sibling, 2 replies; 11+ messages in thread
From: Thomas Schwinge @ 2015-05-20 8:24 UTC (permalink / raw)
To: Cesar Philippidis, Jakub Jelinek; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 11096 bytes --]
Hi!
On Fri, 15 May 2015 11:10:21 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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.
Thanks!
> 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
> ...
> }
OK, but see below.
> 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 preference is to keep these as single files (so that C and C++ can
easily be maintained together), and use the appropriate dg-* directives
to select the expected C or C++ error message, respectively, or use
regular expressions so as to match both the expected C and C++ error
variants in one go, if they're similar enough.
> 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.
What remains to be done?
Then, what about the Fortran front end? Checking already done as well as
test coverage existing, similar to C and C++?
Patch review comments:
> --- 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;
Hmm.
> @@ -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");
> + }
Instead of in c_parser_oacc_shape_clause, shouldn't such checking rather
be done inside the function invoking c_parser_oacc_shape_clause, that is,
c_parser_oacc_parallel, etc.?
> @@ -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);
I would have expected such checking to be done in c_omp_finish_clauses --
But maybe it's also OK to do it here, given that the loop construct is
the only one where these clauses can appear. Jakub, any strong
preference?
> @@ -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;
Regarding this...
> @@ -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;
..., and this: why not do such nesting checking in
gcc/omp-low.c:check_omp_nesting_restrictions? Currently (changed by
Bernd in internal r442824, 2014-11-29) we're allowing all
OpenACC-inside-OpenACC nesting -- shouldn't that be changed instead of
repeating the checks in every front end (Jakub?)?
I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
»gang, worker and vector may occur only once in a loop nest«, and »gang,
worker and vector must occur in this order in a loop nest«. Don't know
if that conceptually also belongs into
gcc/omp-low.c:check_omp_nesting_restrictions?
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
As for C.
> --- a/gcc/cp/parser.h
> +++ b/gcc/cp/parser.h
Likewise.
> --- 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");
> - }
> }
> }
> --- 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 }
I like it better if these are placed next to the line they apply to.
Makes it easier to see what's going on.
> --- /dev/null
> +++ b/gcc/testsuite/g++.dg/goacc/loop-4.C
You didn't actually commit this file.
> @@ -0,0 +1,686 @@
> +// { dg-do compile }
> +// { dg-options "-fopenacc" }
> +// { dg-additional-options "-fmax-errors=200" }
Inside gcc/testsuite/*/goacc/, -fopenacc is enabled by default.
> +
> +int
> +main ()
> +{
> + int i, j;
> +
> +#pragma acc kernels
> + {
> +[...]
> +#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++)
> + {
> +[...]
Doesn't it make sense to separate test cases that are expected to be
handled successfully (no error messages) from those where we expect error
diagnostics, so that we can be sure the former ones will also be
successfully run through all later compiler passes?
> +[...]
> +// { 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 }
Hmm, that doesn't happen for C?
> --- /dev/null
> +++ b/gcc/testsuite/gcc.dg/goacc/loop-1.c
You didn't actually commit this file.
> @@ -0,0 +1,674 @@
> +[...]
> +// { dg-excess-errors "invalid controlling predicate" }
Why is that needed?
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 8:24 ` [patch,gomp4] error on invalid acc loop clauses Thomas Schwinge
@ 2015-05-20 8:56 ` Jakub Jelinek
2015-05-20 9:41 ` Thomas Schwinge
2015-05-20 14:19 ` Cesar Philippidis
1 sibling, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-05-20 8:56 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Cesar Philippidis, gcc-patches
On Wed, May 20, 2015 at 10:23:21AM +0200, Thomas Schwinge wrote:
> > + 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);
>
> I would have expected such checking to be done in c_omp_finish_clauses --
> But maybe it's also OK to do it here, given that the loop construct is
> the only one where these clauses can appear. Jakub, any strong
> preference?
In the C FE, it is kind of arbitrary, some checks are done during parsing
immediately, others are done in c_omp_finish_clauses.
In the C++ FE, obviously more care on where things are diagnosed is needed,
so many more checks are done in finish_omp_clauses, because we might want to
wait until templates are instantiated.
>
> ..., and this: why not do such nesting checking in
> gcc/omp-low.c:check_omp_nesting_restrictions? Currently (changed by
> Bernd in internal r442824, 2014-11-29) we're allowing all
> OpenACC-inside-OpenACC nesting -- shouldn't that be changed instead of
> repeating the checks in every front end (Jakub?)?
Yeah, testing nesting restrictions should be done in omp-low.c if possible.
Adding ugly hacks to the FEs tracking the current state and duplicating
across all 3 FEs is undesirable. Note, in C++ FE we already have sk_omp
so some kind of OpenMP binding scope, but I think we don't have anything
similar in the C FE.
> I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
> »gang, worker and vector may occur only once in a loop nest«, and »gang,
> worker and vector must occur in this order in a loop nest«. Don't know
> if that conceptually also belongs into
> gcc/omp-low.c:check_omp_nesting_restrictions?
Doesn't look like anything related to construct/region nesting...
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 8:56 ` Jakub Jelinek
@ 2015-05-20 9:41 ` Thomas Schwinge
2015-05-20 10:01 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Thomas Schwinge @ 2015-05-20 9:41 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Cesar Philippidis, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1006 bytes --]
Hi!
On Wed, 20 May 2015 10:43:27 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, May 20, 2015 at 10:23:21AM +0200, Thomas Schwinge wrote:
> > I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
> > »gang, worker and vector may occur only once in a loop nest«, and »gang,
> > worker and vector must occur in this order in a loop nest«. Don't know
> > if that conceptually also belongs into
> > gcc/omp-low.c:check_omp_nesting_restrictions?
>
> Doesn't look like anything related to construct/region nesting...
It is checking invalid nesting of loop constructs, for example:
#pragma acc loop gang
for [...]
{
#pragma acc loop gang // gang, worker and vector may occur only once in a loop nest
for [...]
..., or:
#pragma acc loop vector
for [...]
{
#pragma acc loop gang // gang, worker and vector must occur in this order in a loop nest
for [...]
..., and so on.
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 9:41 ` Thomas Schwinge
@ 2015-05-20 10:01 ` Jakub Jelinek
0 siblings, 0 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-05-20 10:01 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Cesar Philippidis, gcc-patches
On Wed, May 20, 2015 at 11:32:20AM +0200, Thomas Schwinge wrote:
> Hi!
>
> On Wed, 20 May 2015 10:43:27 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Wed, May 20, 2015 at 10:23:21AM +0200, Thomas Schwinge wrote:
> > > I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
> > > »gang, worker and vector may occur only once in a loop nest«, and »gang,
> > > worker and vector must occur in this order in a loop nest«. Don't know
> > > if that conceptually also belongs into
> > > gcc/omp-low.c:check_omp_nesting_restrictions?
> >
> > Doesn't look like anything related to construct/region nesting...
>
> It is checking invalid nesting of loop constructs, for example:
>
> #pragma acc loop gang
> for [...]
> {
> #pragma acc loop gang // gang, worker and vector may occur only once in a loop nest
> for [...]
>
> ..., or:
>
> #pragma acc loop vector
> for [...]
> {
> #pragma acc loop gang // gang, worker and vector must occur in this order in a loop nest
> for [...]
>
> ..., and so on.
Ah, in that case it is the right function for that.
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 8:24 ` [patch,gomp4] error on invalid acc loop clauses Thomas Schwinge
2015-05-20 8:56 ` Jakub Jelinek
@ 2015-05-20 14:19 ` Cesar Philippidis
2015-05-20 14:36 ` Jakub Jelinek
1 sibling, 1 reply; 11+ messages in thread
From: Cesar Philippidis @ 2015-05-20 14:19 UTC (permalink / raw)
To: Thomas Schwinge, Jakub Jelinek; +Cc: gcc-patches
On 05/20/2015 01:23 AM, Thomas Schwinge wrote:
>> 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 preference is to keep these as single files (so that C and C++ can
> easily be maintained together), and use the appropriate dg-* directives
> to select the expected C or C++ error message, respectively, or use
> regular expressions so as to match both the expected C and C++ error
> variants in one go, if they're similar enough.
>
>> 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.
>
> What remains to be done?
Jakub made some general comments a couple of weeks ago when you applied
our internal changes to gomp-4_0-branch. I was planning on addressing
those comments first before requesting the merge to trunk.
> Then, what about the Fortran front end? Checking already done as well as
> test coverage existing, similar to C and C++?
Fortran is good.
> Patch review comments:
>
>> --- 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;
>
> Hmm.
What's wrong with this? Fortran does something similar. Besides, this is
only temporary until OpenACC 2.5.
>> @@ -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");
>> + }
>
> Instead of in c_parser_oacc_shape_clause, shouldn't such checking rather
> be done inside the function invoking c_parser_oacc_shape_clause, that is,
> c_parser_oacc_parallel, etc.?
I don't think that will help. c_parser_oacc_shape_clause parses 'gang',
'worker' and 'vector' which aren't available to acc parallel or acc
kernels. Well, they are, be rigged up the front end to split acc
parallel loops and acc kernels loop.
>> @@ -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;
>
> Regarding this...
>
>> @@ -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;
>
> ..., and this: why not do such nesting checking in
> gcc/omp-low.c:check_omp_nesting_restrictions? Currently (changed by
> Bernd in internal r442824, 2014-11-29) we're allowing all
> OpenACC-inside-OpenACC nesting -- shouldn't that be changed instead of
> repeating the checks in every front end (Jakub?)?
The fortran front end is doing this. Also, Joseph told me the front ends
should report error messages when possible. I have no problems reverting
back to the original behavior though.
> I see that some checking is also being done gcc/omp-low.c:scan_omp_for:
> »gang, worker and vector may occur only once in a loop nest«, and »gang,
> worker and vector must occur in this order in a loop nest«. Don't know
> if that conceptually also belongs into
> gcc/omp-low.c:check_omp_nesting_restrictions?
Yeah, someone needs to clean that up. I tried to keep this patch local
to the c and c++ front ends.
>> --- 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 }
>
> I like it better if these are placed next to the line they apply to.
> Makes it easier to see what's going on.
There's already a dg-error on those lines and I think dg-error can only
report one error per line.
>> --- /dev/null
>> +++ b/gcc/testsuite/g++.dg/goacc/loop-4.C
>
> You didn't actually commit this file.
>
>> @@ -0,0 +1,686 @@
>> +// { dg-do compile }
>> +// { dg-options "-fopenacc" }
>> +// { dg-additional-options "-fmax-errors=200" }
>
> Inside gcc/testsuite/*/goacc/, -fopenacc is enabled by default.
True. However the toplevel dg.exp also executes those tests and it
doesn't use -fopenacc without dg-options.
>> +
>> +int
>> +main ()
>> +{
>> + int i, j;
>> +
>> +#pragma acc kernels
>> + {
>> +[...]
>> +#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++)
>> + {
>> +[...]
>
> Doesn't it make sense to separate test cases that are expected to be
> handled successfully (no error messages) from those where we expect error
> diagnostics, so that we can be sure the former ones will also be
> successfully run through all later compiler passes?
I thought the ones that we expect to run successfully belong in libgomp?
>> +[...]
>
>> +// { 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 }
>
> Hmm, that doesn't happen for C?
I'm not sure why, unfortunately.
>> --- /dev/null
>> +++ b/gcc/testsuite/gcc.dg/goacc/loop-1.c
>
> You didn't actually commit this file.
Two files that I missed.
>> @@ -0,0 +1,674 @@
>> +[...]
>> +// { dg-excess-errors "invalid controlling predicate" }
>
> Why is that needed?
I'm not sure. I thought it had something to do with Bernd's predication
pass.
Cesar
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 14:19 ` Cesar Philippidis
@ 2015-05-20 14:36 ` Jakub Jelinek
2015-05-20 23:35 ` Cesar Philippidis
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-05-20 14:36 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: Thomas Schwinge, gcc-patches
On Wed, May 20, 2015 at 07:17:35AM -0700, Cesar Philippidis wrote:
> > ..., and this: why not do such nesting checking in
> > gcc/omp-low.c:check_omp_nesting_restrictions? Currently (changed by
> > Bernd in internal r442824, 2014-11-29) we're allowing all
> > OpenACC-inside-OpenACC nesting -- shouldn't that be changed instead of
> > repeating the checks in every front end (Jakub?)?
>
> The fortran front end is doing this. Also, Joseph told me the front ends
> should report error messages when possible. I have no problems reverting
> back to the original behavior though.
For OpenMP/OpenACC, there is still lots of diagnostics emitted during
gimplification and at the start of the omp lowering phases, so for
diagnostics purposes you can consider them as part of a common layer for all
the 3 FEs.
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 14:36 ` Jakub Jelinek
@ 2015-05-20 23:35 ` Cesar Philippidis
2015-06-03 0:03 ` [ping] " Cesar Philippidis
0 siblings, 1 reply; 11+ messages in thread
From: Cesar Philippidis @ 2015-05-20 23:35 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Thomas Schwinge, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 534 bytes --]
On 05/20/2015 07:32 AM, Jakub Jelinek wrote:
> For OpenMP/OpenACC, there is still lots of diagnostics emitted during
> gimplification and at the start of the omp lowering phases, so for
> diagnostics purposes you can consider them as part of a common layer for all
> the 3 FEs.
Thanks, that makes sense. I didn't touch the fortran front end, but I
did revert most of my c and c++ front end changes in gomp-4_0-branch to
defer all of the common error checking to the omp lowering phase.
Is this patch ok for gomp-4_0-branch?
Cesar
[-- Attachment #2: gwv-loops-cleanup.diff --]
[-- Type: text/x-patch, Size: 34490 bytes --]
2015-05-20 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-parser.c (typedef struct c_parser): Remove oacc_parallel_region
and oacc_kernels_region.
(c_parser_oacc_shape_clause): Don't check for parallel-specific
clauses here.
(c_parser_oacc_loop): Don't check for incompatible clauses.
(c_parser_oacc_kernels): Don't check for nested parallelism.
(c_parser_oacc_parallel): Likewise.
gcc/cp/
* parser.h: Remove oacc_parallel_region and oacc_kernels_region.
* parser.c (cp_parser_oacc_shape_clause): Don't check for
parallel-specific clauses here.
(cp_parser_oacc_loop): Don't check for incompatible clauses.
(cp_parser_oacc_parallel_kernels): Don't check for nested parallelism.
gcc/
* omp-low.c (enclosing_target_ctx): Un-ifdef. Remove null checking
assert.
(scan_omp_for): Check for incompatible combination of acc loop clauses.
gcc/testsuite/
* c-c++-common/goacc/loop-2.c (main): New test.
* c-c++-common/goacc/nesting-fail-1.c (f_acc_parallel): Update error
messages.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index f508b91..4099cc4 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -234,10 +234,6 @@ 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. */
@@ -10857,17 +10853,6 @@ 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;
}
@@ -12737,10 +12722,7 @@ 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, c;
- bool gwv = false;
- bool auto_clause = false;
- bool seq_clause = false;
+ tree stmt, clauses, block;
strcat (p_name, " loop");
mask |= OACC_LOOP_CLAUSE_MASK;
@@ -12751,33 +12733,6 @@ 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);
@@ -12820,13 +12775,6 @@ 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))
{
@@ -12840,7 +12788,6 @@ 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;
}
}
@@ -12851,7 +12798,6 @@ 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;
}
@@ -12898,13 +12844,6 @@ 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))
@@ -12919,7 +12858,6 @@ 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;
}
}
@@ -12929,7 +12867,6 @@ 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 2947bf4..748905c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -28335,17 +28335,6 @@ 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;
}
@@ -32210,10 +32199,7 @@ 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, c;
- bool gwv = false;
- bool auto_clause = false;
- bool seq_clause = false;
+ tree stmt, clauses, block;
int save;
strcat (p_name, " loop");
@@ -32226,33 +32212,6 @@ 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);
@@ -32330,21 +32289,13 @@ 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");
}
@@ -32363,10 +32314,6 @@ 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;
}
}
@@ -32383,10 +32330,6 @@ 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 07292bf..8eb5484 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -377,11 +377,6 @@ 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 3414ab5..67ca302 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2892,18 +2892,14 @@ finish_taskreg_scan (omp_context *ctx)
}
}
-
-#if 0
static omp_context *
enclosing_target_ctx (omp_context *ctx)
{
while (ctx != NULL
&& gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
ctx = ctx->outer;
- gcc_assert (ctx != NULL);
return ctx;
}
-#endif
static bool
oacc_loop_or_target_p (gimple stmt)
@@ -2927,6 +2923,9 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
omp_context *ctx;
size_t i;
tree clauses = gimple_omp_for_clauses (stmt);
+ bool gwv_clause = false;
+ bool auto_clause = false;
+ bool seq_clause = false;
if (outer_ctx)
outer_type = gimple_code (outer_ctx->stmt);
@@ -2941,11 +2940,30 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
{
int val;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
- val = MASK_GANG;
+ {
+ val = MASK_GANG;
+ gwv_clause = true;
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)
- val = MASK_WORKER;
+ {
+ val = MASK_WORKER;
+ gwv_clause = true;
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)
- val = MASK_VECTOR;
+ {
+ val = MASK_VECTOR;
+ gwv_clause = true;
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SEQ)
+ {
+ seq_clause = true;
+ continue;
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AUTO)
+ {
+ auto_clause = true;
+ continue;
+ }
else
continue;
ctx->gwv_this |= val;
@@ -2961,9 +2979,24 @@ 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);
+ /* Enclosing may be null if we are inside an acc routine. If
+ that's the case, treat this loop as a parallel. */
+ if (enclosing == NULL || 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");
+ }
}
}
+ if ((gwv_clause && auto_clause) || (auto_clause && seq_clause))
+ error_at (gimple_location (stmt), "incompatible use of clause auto");
+ else if (gwv_clause && seq_clause)
+ error_at (gimple_location (stmt), "incompatible use of clause seq");
+
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2.c b/gcc/testsuite/c-c++-common/goacc/loop-2.c
new file mode 100644
index 0000000..7265614
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2.c
@@ -0,0 +1,673 @@
+/* { 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected" }
+ 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 "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5) // { dg-error "no arguments allowed to gang" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5) // { dg-error "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5) // { dg-error "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5) // { dg-error "no arguments allowed to gang" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected" }
+ 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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq gang // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc kernels loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile() // { dg-error "expected" }
+ 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 "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop gang(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq gang // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector(length:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop vector // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc parallel loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile() // { dg-error "expected" }
+ 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;
+}
+
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 6054fb8..9bc95e9 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -7,7 +7,7 @@ f_acc_parallel (void)
{
#pragma acc parallel
{
-#pragma acc parallel /* { dg-error "nested parallel region" } */
+#pragma acc parallel /* { dg-error "parallel construct inside of parallel region" } */
;
#pragma acc kernels /* { dg-error " kernels construct inside of parallel" } */
;
@@ -26,9 +26,9 @@ f_acc_kernels (void)
{
#pragma acc kernels
{
-#pragma acc parallel /* { dg-error "nested parallel region" } */
+#pragma acc parallel /* { dg-error "parallel construct inside of kernels region" } */
;
-#pragma acc kernels /* { dg-error "nested kernels region" } */
+#pragma acc kernels /* { dg-error "kernels construct inside of kernels region" } */
;
#pragma acc data /* { dg-error "data construct inside of kernels region" } */
;
@@ -37,8 +37,3 @@ 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 }
^ permalink raw reply [flat|nested] 11+ messages in thread
* [ping] Re: [patch,gomp4] error on invalid acc loop clauses
2015-05-20 23:35 ` Cesar Philippidis
@ 2015-06-03 0:03 ` Cesar Philippidis
0 siblings, 0 replies; 11+ messages in thread
From: Cesar Philippidis @ 2015-06-03 0:03 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Thomas Schwinge, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 719 bytes --]
Any thoughts? Maybe I should just apply this patch to gomp-4_0-branch
since it doesn't introduce any regressions.
On 05/20/2015 04:26 PM, Cesar Philippidis wrote:
> On 05/20/2015 07:32 AM, Jakub Jelinek wrote:
>
>> For OpenMP/OpenACC, there is still lots of diagnostics emitted during
>> gimplification and at the start of the omp lowering phases, so for
>> diagnostics purposes you can consider them as part of a common layer for all
>> the 3 FEs.
>
> Thanks, that makes sense. I didn't touch the fortran front end, but I
> did revert most of my c and c++ front end changes in gomp-4_0-branch to
> defer all of the common error checking to the omp lowering phase.
>
> Is this patch ok for gomp-4_0-branch?
Cesar
[-- Attachment #2: gwv-loops-cleanup.diff --]
[-- Type: text/x-patch, Size: 34490 bytes --]
2015-05-20 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-parser.c (typedef struct c_parser): Remove oacc_parallel_region
and oacc_kernels_region.
(c_parser_oacc_shape_clause): Don't check for parallel-specific
clauses here.
(c_parser_oacc_loop): Don't check for incompatible clauses.
(c_parser_oacc_kernels): Don't check for nested parallelism.
(c_parser_oacc_parallel): Likewise.
gcc/cp/
* parser.h: Remove oacc_parallel_region and oacc_kernels_region.
* parser.c (cp_parser_oacc_shape_clause): Don't check for
parallel-specific clauses here.
(cp_parser_oacc_loop): Don't check for incompatible clauses.
(cp_parser_oacc_parallel_kernels): Don't check for nested parallelism.
gcc/
* omp-low.c (enclosing_target_ctx): Un-ifdef. Remove null checking
assert.
(scan_omp_for): Check for incompatible combination of acc loop clauses.
gcc/testsuite/
* c-c++-common/goacc/loop-2.c (main): New test.
* c-c++-common/goacc/nesting-fail-1.c (f_acc_parallel): Update error
messages.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index f508b91..4099cc4 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -234,10 +234,6 @@ 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. */
@@ -10857,17 +10853,6 @@ 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;
}
@@ -12737,10 +12722,7 @@ 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, c;
- bool gwv = false;
- bool auto_clause = false;
- bool seq_clause = false;
+ tree stmt, clauses, block;
strcat (p_name, " loop");
mask |= OACC_LOOP_CLAUSE_MASK;
@@ -12751,33 +12733,6 @@ 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);
@@ -12820,13 +12775,6 @@ 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))
{
@@ -12840,7 +12788,6 @@ 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;
}
}
@@ -12851,7 +12798,6 @@ 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;
}
@@ -12898,13 +12844,6 @@ 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))
@@ -12919,7 +12858,6 @@ 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;
}
}
@@ -12929,7 +12867,6 @@ 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 2947bf4..748905c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -28335,17 +28335,6 @@ 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;
}
@@ -32210,10 +32199,7 @@ 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, c;
- bool gwv = false;
- bool auto_clause = false;
- bool seq_clause = false;
+ tree stmt, clauses, block;
int save;
strcat (p_name, " loop");
@@ -32226,33 +32212,6 @@ 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);
@@ -32330,21 +32289,13 @@ 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");
}
@@ -32363,10 +32314,6 @@ 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;
}
}
@@ -32383,10 +32330,6 @@ 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 07292bf..8eb5484 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -377,11 +377,6 @@ 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 3414ab5..67ca302 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2892,18 +2892,14 @@ finish_taskreg_scan (omp_context *ctx)
}
}
-
-#if 0
static omp_context *
enclosing_target_ctx (omp_context *ctx)
{
while (ctx != NULL
&& gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
ctx = ctx->outer;
- gcc_assert (ctx != NULL);
return ctx;
}
-#endif
static bool
oacc_loop_or_target_p (gimple stmt)
@@ -2927,6 +2923,9 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
omp_context *ctx;
size_t i;
tree clauses = gimple_omp_for_clauses (stmt);
+ bool gwv_clause = false;
+ bool auto_clause = false;
+ bool seq_clause = false;
if (outer_ctx)
outer_type = gimple_code (outer_ctx->stmt);
@@ -2941,11 +2940,30 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
{
int val;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
- val = MASK_GANG;
+ {
+ val = MASK_GANG;
+ gwv_clause = true;
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)
- val = MASK_WORKER;
+ {
+ val = MASK_WORKER;
+ gwv_clause = true;
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)
- val = MASK_VECTOR;
+ {
+ val = MASK_VECTOR;
+ gwv_clause = true;
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SEQ)
+ {
+ seq_clause = true;
+ continue;
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AUTO)
+ {
+ auto_clause = true;
+ continue;
+ }
else
continue;
ctx->gwv_this |= val;
@@ -2961,9 +2979,24 @@ 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);
+ /* Enclosing may be null if we are inside an acc routine. If
+ that's the case, treat this loop as a parallel. */
+ if (enclosing == NULL || 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");
+ }
}
}
+ if ((gwv_clause && auto_clause) || (auto_clause && seq_clause))
+ error_at (gimple_location (stmt), "incompatible use of clause auto");
+ else if (gwv_clause && seq_clause)
+ error_at (gimple_location (stmt), "incompatible use of clause seq");
+
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2.c b/gcc/testsuite/c-c++-common/goacc/loop-2.c
new file mode 100644
index 0000000..7265614
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2.c
@@ -0,0 +1,673 @@
+/* { 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected" }
+ 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 "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang(num:5) // { dg-error "no arguments allowed to gang" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(5) // { dg-error "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker(num:5) // { dg-error "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 "no arguments allowed to gang" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector(length:5) // { dg-error "no arguments allowed to gang" }
+ 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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#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 // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+ for (i = 0; i < 10; i++)
+ { }
+
+#pragma acc loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc loop tile() // { dg-error "expected" }
+ 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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq gang // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc kernels loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc kernels loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc kernels loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc kernels loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc kernels loop tile() // { dg-error "expected" }
+ 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 "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop gang(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#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 "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq gang // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop worker(5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector(length:5) // { dg-error "no arguments allowed to gang" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector
+ for (i = 0; i < 10; i++)
+ {
+#pragma acc parallel loop vector // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop worker // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+#pragma acc parallel loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
+ for (j = 1; j < 10; j++)
+ { }
+ }
+#pragma acc parallel loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#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 // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+#pragma acc parallel loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+ { }
+
+#pragma acc parallel loop tile // { dg-error "expected" }
+ for (i = 0; i < 10; i++)
+ { }
+#pragma acc parallel loop tile() // { dg-error "expected" }
+ 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;
+}
+
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 6054fb8..9bc95e9 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -7,7 +7,7 @@ f_acc_parallel (void)
{
#pragma acc parallel
{
-#pragma acc parallel /* { dg-error "nested parallel region" } */
+#pragma acc parallel /* { dg-error "parallel construct inside of parallel region" } */
;
#pragma acc kernels /* { dg-error " kernels construct inside of parallel" } */
;
@@ -26,9 +26,9 @@ f_acc_kernels (void)
{
#pragma acc kernels
{
-#pragma acc parallel /* { dg-error "nested parallel region" } */
+#pragma acc parallel /* { dg-error "parallel construct inside of kernels region" } */
;
-#pragma acc kernels /* { dg-error "nested kernels region" } */
+#pragma acc kernels /* { dg-error "kernels construct inside of kernels region" } */
;
#pragma acc data /* { dg-error "data construct inside of kernels region" } */
;
@@ -37,8 +37,3 @@ 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 }
^ permalink raw reply [flat|nested] 11+ messages in thread
end of thread, other threads:[~2015-06-02 22:54 UTC | newest]
Thread overview: 11+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <555A05BA.3050001@mentor.com>
2015-05-19 8:04 ` [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used Tom de Vries
2015-05-15 18:14 ` [patch,gomp4] error on invalid acc loop clauses Cesar Philippidis
2015-05-20 7:48 ` [gomp4] bootstrap broken, function enclosing_target_ctx defined but not used Thomas Schwinge
2015-05-20 8:24 ` [patch,gomp4] error on invalid acc loop clauses Thomas Schwinge
2015-05-20 8:56 ` Jakub Jelinek
2015-05-20 9:41 ` Thomas Schwinge
2015-05-20 10:01 ` Jakub Jelinek
2015-05-20 14:19 ` Cesar Philippidis
2015-05-20 14:36 ` Jakub Jelinek
2015-05-20 23:35 ` Cesar Philippidis
2015-06-03 0:03 ` [ping] " Cesar Philippidis
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).