From patchwork Tue Oct 2 14:06:12 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 977861 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-486787-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="qDOI3ZId"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 42Pgtf17gvz9vDh for ; Wed, 3 Oct 2018 00:06:31 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; q=dns; s= default; b=IFVZGYylo8NQUeeA1A1Cn/+ATRvtxuz9VEapar/aj/U0utq32nFCF oMN/mGR7JVNIS5cWZSgwBAilWoiL2e0hIuXa4VVWJwtyLm64f0yh1z+Og79pNX3k gky9P9Xdcxw6dUTPzvj/2Gwb0828tn3ZPffWjzgmNQj84GC+mxAa3M= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:message-id:date:mime-version:content-type; s= default; bh=sZ5hUhG55er/YpR1QT1u7795S3Q=; b=qDOI3ZIde5u/BI/rH+q4 03DkENxtftjPjtckN0YIsL+XvGtuAx8Urv7KZrFfdTu8cY17qKNYQUuhLLwE4c7Q oi1opTjMEnnK/MGB3iUjA9BjRyEyopXAPm+hxZEkXHvusNoQXUqp3P2SY89zSvZC u2j8Px75Wx7wfZRwzhOLrrg= Received: (qmail 75353 invoked by alias); 2 Oct 2018 14:06:23 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 75342 invoked by uid 89); 2 Oct 2018 14:06:22 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=introduces, gang, verifying, gcc_unreachable 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; Tue, 02 Oct 2018 14:06:17 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1g7LJX-0000zX-Iz from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Tue, 02 Oct 2018 07:06:15 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 2 Oct 2018 07:06:12 -0700 From: Cesar Philippidis Subject: [patch,openacc] Use oacc_verify_routine_clauses for C/C++ To: "gcc-patches@gcc.gnu.org" , Thomas Schwinge Message-ID: Date: Tue, 2 Oct 2018 07:06:12 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.9.1 MIME-Version: 1.0 This patch introduces a new oacc_verify_routine_clauses function that reports errors if the user abuses the gang, worker and vector clauses for acc routine directives in C/C++. Fortran is a little different, because the FE has it's own IR. So, while it would be possible to defer checking for gang, worker, vector parallelism until a tree node is created for a function, we'd still have problems of verifying the parallelism for functions and subroutines defined and declared inside modules. The C and C++ FE's are similar enough were they can share a common function. Is this OK for trunk? I bootstrapped and regression tested it for x86_64 Linux with nvptx offloading. This is only touches the OpenACC code path. Cesar [OpenACC] Use oacc_verify_routine_clauses for C/C++ 2018-XX-YY Thomas Schwinge Cesar Philippidis gcc/ * omp-general.c (oacc_build_routine_dims): Move some of its processing into... (oacc_verify_routine_clauses): ... this new function. * omp-general.h (oacc_verify_routine_clauses): New prototype. gcc/c/ * c-parser.c (c_parser_oacc_routine): Normalize order of clauses. (c_finish_oacc_routine): Call oacc_verify_routine_clauses. gcc/cp/ * parser.c (cp_parser_oacc_routine) (cp_parser_late_parsing_oacc_routine): Normalize order of clauses. (cp_finalize_oacc_routine): Call oacc_verify_routine_clauses. gcc/testsuite/ * c-c++-common/goacc/routine-level-of-parallelism-1.c: New test. (cherry picked from gomp-4_0-branch r239520) --- gcc/c/c-parser.c | 8 + gcc/cp/parser.c | 9 + gcc/omp-general.c | 69 ++++- gcc/omp-general.h | 1 + .../goacc/routine-level-of-parallelism-1.c | 265 ++++++++++++++++++ 5 files changed, 342 insertions(+), 10 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 3ca8fe71cc4..3517cb783d9 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -14999,6 +14999,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context) data.clauses = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); if (TREE_CODE (decl) != FUNCTION_DECL) { @@ -15013,6 +15016,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context) data.clauses = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); /* Emit a helpful diagnostic if there's another pragma following this one. Also don't allow a static assertion declaration, as in the @@ -15076,6 +15082,8 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, return; } + oacc_verify_routine_clauses (&data->clauses, data->loc); + if (oacc_get_fn_attrib (fndecl)) { error_at (data->loc, diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 241226d8c21..fa7ee7798ae 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -38117,6 +38117,9 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine", cp_lexer_peek_token (parser->lexer)); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); if (decl && is_overloaded_fn (decl) && (TREE_CODE (decl) != FUNCTION_DECL @@ -38213,6 +38216,9 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) parser->oacc_routine->clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine", pragma_tok); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses); cp_parser_pop_lexer (parser); /* Later, cp_finalize_oacc_routine will process the clauses, and then set fndecl_seen. */ @@ -38247,6 +38253,9 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) return; } + oacc_verify_routine_clauses (&parser->oacc_routine->clauses, + parser->oacc_routine->loc); + if (oacc_get_fn_attrib (fndecl)) { error_at (parser->oacc_routine->loc, diff --git a/gcc/omp-general.c b/gcc/omp-general.c index cabbbbc6de2..3ea2224957d 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -559,9 +559,64 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec *args) } } -/* Process the routine's dimension clauess to generate an attribute - value. Issue diagnostics as appropriate. We default to SEQ - (OpenACC 2.5 clarifies this). All dimensions have a size of zero +/* Verify OpenACC routine clauses. + + The chain of clauses returned will contain exactly one clause specifying the + level of parallelism. */ + +void +oacc_verify_routine_clauses (tree *clauses, location_t loc) +{ + tree c_level = NULL_TREE; + tree c_p = NULL_TREE; + for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_SEQ: + if (c_level == NULL_TREE) + c_level = c; + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level)) + { + /* This has already been diagnosed in the front ends. */ + /* Drop the duplicate clause. */ + gcc_checking_assert (c_p != NULL_TREE); + OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c); + c = c_p; + } + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qs specifies a conflicting level of parallelism", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + inform (OMP_CLAUSE_LOCATION (c_level), + "... to the previous %qs clause here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]); + /* Drop the conflicting clause. */ + gcc_checking_assert (c_p != NULL_TREE); + OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c); + c = c_p; + } + break; + case OMP_CLAUSE_NOHOST: + break; + default: + gcc_unreachable (); + } + if (c_level == NULL_TREE) + { + /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a + implementation add an implicit "seq" clause. */ + c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ); + OMP_CLAUSE_CHAIN (c_level) = *clauses; + *clauses = c_level; + } +} + +/* Process the OpenACC routine's clauses to generate an attribute + for the level of parallelism. All dimensions have a size of zero (dynamic). TREE_PURPOSE is set to indicate whether that dimension can have a loop partitioned on it. non-zero indicates yes, zero indicates no. By construction once a non-zero has been @@ -583,16 +638,10 @@ oacc_build_routine_dims (tree clauses) for (ix = GOMP_DIM_MAX + 1; ix--;) if (OMP_CLAUSE_CODE (clauses) == ids[ix]) { - if (level >= 0) - error_at (OMP_CLAUSE_LOCATION (clauses), - "multiple loop axes specified for routine"); level = ix; break; } - - /* Default to SEQ. */ - if (level < 0) - level = GOMP_DIM_MAX; + gcc_checking_assert (level >= 0); tree dims = NULL_TREE; diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 66f0a33c2e2..e0de288dd9a 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -83,6 +83,7 @@ extern int omp_max_simt_vf (void); extern tree oacc_launch_pack (unsigned code, tree device, unsigned op); extern void oacc_replace_fn_attrib (tree fn, tree dims); extern void oacc_set_fn_attrib (tree fn, tree clauses, vec *args); +extern void oacc_verify_routine_clauses (tree *, location_t); extern tree oacc_build_routine_dims (tree clauses); extern tree oacc_get_fn_attrib (tree fn); extern bool offloading_function_p (tree fn); diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c new file mode 100644 index 00000000000..d71d6e088f7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c @@ -0,0 +1,265 @@ +/* Test various aspects of clauses specifying incompatible levels of + parallelism with the OpenACC routine directive. The Fortran counterpart is + ../../gfortran.dg/goacc/routine-level-of-parallelism-1.f. */ + +extern void g_1 (void); +#pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */ + +#pragma acc routine worker worker /* { dg-error "too many 'worker' clauses" } */ +void w_1 (void) +{ +} + +#pragma acc routine vector vector /* { dg-error "too many 'vector' clauses" } */ +void v_1 (void) +{ +} + +#pragma acc routine seq seq /* { dg-error "too many 'seq' clauses" } */ +extern void s_1 (void); + + +#pragma acc routine gang gang gang /* { dg-error "too many 'gang' clauses" } */ +void g_2 (void) +{ +} + +#pragma acc routine worker worker worker /* { dg-error "too many 'worker' clauses" } */ +extern void w_2 (void); + +extern void v_2 (void); +#pragma acc routine (v_2) vector vector vector /* { dg-error "too many 'vector' clauses" } */ + +#pragma acc routine seq seq seq /* { dg-error "too many 'seq' clauses" } */ +void s_2 (void) +{ +} + + +#pragma acc routine \ + gang \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ +void g_3 (void) +{ +} +#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \ + gang \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ +#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \ + gang \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ + +extern void w_3 (void); +#pragma acc routine (w_3) \ + worker \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ +#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \ + worker \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \ + worker \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ + +#pragma acc routine \ + vector \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ +void v_3 (void) +{ +} +#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \ + vector \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ +#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \ + vector \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ + +extern void s_3 (void); +#pragma acc routine (s_3) \ + seq \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \ + seq \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ +#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \ + seq \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ + + +#pragma acc routine \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ +extern void g_4 (void); +#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ +#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ + +extern void w_4 (void); +#pragma acc routine (w_4) \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ + +#pragma acc routine \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +void v_4 (void) +{ +} +#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ + +#pragma acc routine \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +void s_4 (void) +{ +} +#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ +#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ + + +#pragma acc routine \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +void g_5 (void) +{ +} +#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ + +#pragma acc routine \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +extern void w_5 (void); +#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ + +#pragma acc routine \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +extern void v_5 (void); +#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ + +extern void s_5 (void); +#pragma acc routine (s_5) \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ -- 2.17.1