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.
@@ -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;
}
@@ -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;
}
@@ -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;
@@ -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);
new file mode 100644
@@ -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;
+}
+
@@ -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 }