From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 4411 invoked by alias); 20 May 2015 23:27:04 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 4399 invoked by uid 89); 20 May 2015 23:27:02 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.6 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_LOW,SPF_PASS,UNWANTED_LANGUAGE_BODY autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 20 May 2015 23:26:58 +0000 Received: from svr-orw-fem-04.mgc.mentorg.com ([147.34.97.41]) by relay1.mentorg.com with esmtp id 1YvDO2-00000Y-G1 from Cesar_Philippidis@mentor.com ; Wed, 20 May 2015 16:26:54 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-04.mgc.mentorg.com (147.34.97.41) with Microsoft SMTP Server id 14.3.224.2; Wed, 20 May 2015 16:26:54 -0700 Message-ID: <555D183D.5090406@codesourcery.com> Date: Wed, 20 May 2015 23:35:00 -0000 From: Cesar Philippidis User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.6.0 MIME-Version: 1.0 To: Jakub Jelinek CC: Thomas Schwinge , "gcc-patches@gcc.gnu.org" Subject: Re: [patch,gomp4] error on invalid acc loop clauses References: <5556368D.7010904@codesourcery.com> <87mw0zirnq.fsf@schwinge.name> <555C977F.4000001@codesourcery.com> <20150520143255.GJ1751@tucnak.redhat.com> In-Reply-To: <20150520143255.GJ1751@tucnak.redhat.com> Content-Type: multipart/mixed; boundary="------------020302020107080002050903" X-SW-Source: 2015-05/txt/msg01899.txt.bz2 --------------020302020107080002050903 Content-Type: text/plain; charset="windows-1252" Content-Transfer-Encoding: 7bit Content-length: 534 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 --------------020302020107080002050903 Content-Type: text/x-patch; name="gwv-loops-cleanup.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="gwv-loops-cleanup.diff" Content-length: 34490 2015-05-20 Cesar Philippidis 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 %"); - else if (gwv && seq_clause) - c_parser_error (parser, "incompatible use of clause %"); - else if (auto_clause && seq_clause) - c_parser_error (parser, "incompatible use of clause % and %"); - 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 %"); - else if (gwv && seq_clause) - cp_parser_error (parser, "incompatible use of clause %"); - else if (auto_clause && seq_clause) - cp_parser_error (parser, "incompatible use of clause % and %"); - 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 *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 } --------------020302020107080002050903--