From patchwork Wed Aug 17 01:05:21 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 659865 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 3sDWHz0xWhz9t1F for ; Wed, 17 Aug 2016 11:05:48 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=Md0PkBUo; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type:content-transfer-encoding; q=dns; s= default; b=HqClRxFah2Ugn8O8hJmQ+5P4ycAnxAJAItDr1AgW3gZiOUiCBCUw1 oZ1Sz4uJ2BK43lNGXSUAVl1n3gjoFNUYE1gOFg8Bzeerhv+Kko7XyESbBuqG5lro EIU1ZcLuE5RGsXbFVKXsVJAfOLoK4dUAdI3vw71vm5qryK0GSX/Jpo= 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 :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type:content-transfer-encoding; s=default; bh=ialeJcvbdwG36vbdZCtq4OYThlE=; b=Md0PkBUoAgUZMCaZwVogCjjYtBl1 xafZhpNn5pzTp/Mp7pNJC3bSrIEQ0DbiN52eaSxc0dbLkOEX7RuJh4Xd9fS9Y8Ak CH2BvWW+pHfNy/iRXyuP/TaBRdd9w1NIQrt2GYLSimC13irFeh/i88Ufthqh5ZhH sK/hDd6olsfsxtg= Received: (qmail 20360 invoked by alias); 17 Aug 2016 01:05:40 -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 20168 invoked by uid 89); 17 Aug 2016 01:05:38 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.2 required=5.0 tests=AWL, BAYES_20, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=ordinary, Collect, NOTHING, THREE 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, 17 Aug 2016 01:05:27 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1bZpIK-0003rW-Dm from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Tue, 16 Aug 2016 18:05:25 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Wed, 17 Aug 2016 02:05:22 +0100 From: Thomas Schwinge To: CC: Subject: Re: Repeated use of the OpenACC routine directive In-Reply-To: <87vazkij2b.fsf@hertz.schwinge.homeip.net> References: <87vazkij2b.fsf@hertz.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Wed, 17 Aug 2016 03:05:21 +0200 Message-ID: <877fbggq7i.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Mon, 01 Aug 2016 17:51:24 +0200, I wrote: > We found that it's not correct that we currently unconditionally diagnose > an error for repeated use of the OpenACC routine directive on one > function/declaration. (For reference, it is also permissible for an > "ordinary" function to have several declarations plus a definition, as > long as these are compatible.) This is, the following shall be valid: > > #pragma acc routine worker > void f(void) > { > } > #pragma acc routine (f) worker > #pragma acc routine worker > extern void f(void); > > [...] In r239521 committed to gomp-4_0-branch: commit bffb0ee6c0a83b8c85cd919e1172086b51fdc452 Author: tschwinge Date: Wed Aug 17 00:55:02 2016 +0000 Repeated use of the C/C++ OpenACC routine directive gcc/ * omp-low.c (verify_oacc_routine_clauses): Change formal parameters. Add checking if already marked as an accelerator routine. Adjust all users. gcc/c/ * c-parser.c (c_finish_oacc_routine): Rework checking if already marked as an accelerator routine. gcc/cp/ * parser.c (cp_finalize_oacc_routine): Rework checking if already marked as an accelerator routine. gcc/testsuite/ * c-c++-common/goacc/oaccdevlow-routine.c: Update. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-2.c: New file. gcc/testsuite/ * c-c++-common/goacc/routine-1.c: Update. * c-c++-common/goacc/routine-2.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * g++.dg/goacc/routine-2.C: Likewise. * c-c++-common/goacc/routine-nohost-2.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@239521 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 4 + gcc/c/ChangeLog.gomp | 3 + gcc/c/c-parser.c | 42 ++-- gcc/cp/ChangeLog.gomp | 3 + gcc/cp/parser.c | 51 +++-- gcc/omp-low.c | 146 ++++++++++++- gcc/omp-low.h | 2 +- gcc/testsuite/ChangeLog.gomp | 11 + gcc/testsuite/c-c++-common/goacc/routine-1.c | 51 ++++- gcc/testsuite/c-c++-common/goacc/routine-2.c | 143 +++++++++++++ gcc/testsuite/c-c++-common/goacc/routine-5.c | 46 +--- .../goacc/routine-level-of-parallelism-1.c | 233 ++++++++++++++++++--- .../goacc/routine-level-of-parallelism-2.c | 73 +++++++ .../c-c++-common/goacc/routine-nohost-1.c | 20 ++ .../c-c++-common/goacc/routine-nohost-2.c | 97 +++++++++ gcc/testsuite/g++.dg/goacc/routine-2.C | 9 + 16 files changed, 816 insertions(+), 118 deletions(-) Grüße Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index bcd011d..f2e8f0d 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,9 @@ 2016-08-17 Thomas Schwinge + * omp-low.c (verify_oacc_routine_clauses): Change formal + parameters. Add checking if already marked as an accelerator + routine. Adjust all users. + * omp-low.c (build_oacc_routine_dims): Move some of its processing into... (verify_oacc_routine_clauses): ... this new function. diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 37f5602..57b8dc9 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,5 +1,8 @@ 2016-08-17 Thomas Schwinge + * c-parser.c (c_finish_oacc_routine): Rework checking if already + marked as an accelerator routine. + * c-parser.c (c_parser_oacc_routine): Normalize order of clauses. (c_finish_oacc_routine): Call verify_oacc_routine_clauses. diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 3d50676..824ee14 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -14289,33 +14289,37 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, return; } - verify_oacc_routine_clauses (&data->clauses, data->loc); - - if (get_oacc_fn_attrib (fndecl)) + int compatible + = verify_oacc_routine_clauses (fndecl, &data->clauses, data->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (data->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); data->error_seen = true; return; } - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + if (compatible > 0) { - error_at (data->loc, - "%<#pragma acc routine%> must be applied before %s", - TREE_USED (fndecl) ? "use" : "definition"); - data->error_seen = true; - return; } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (data->loc, + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + data->error_seen = true; + return; + } - /* Process the routine's dimension clauses. */ - tree dims = build_oacc_routine_dims (data->clauses); - replace_oacc_fn_attrib (fndecl, dims); + /* Set the routine's level of parallelism. */ + tree dims = build_oacc_routine_dims (data->clauses); + replace_oacc_fn_attrib (fndecl, dims); - /* Add an "omp declare target" attribute. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - data->clauses, DECL_ATTRIBUTES (fndecl)); + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + data->clauses, DECL_ATTRIBUTES (fndecl)); + } /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp index 7f634d9..bdb4502 100644 --- gcc/cp/ChangeLog.gomp +++ gcc/cp/ChangeLog.gomp @@ -1,5 +1,8 @@ 2016-08-17 Thomas Schwinge + * parser.c (cp_finalize_oacc_routine): Rework checking if already + marked as an accelerator routine. + * parser.c (cp_parser_oacc_routine) (cp_parser_late_parsing_oacc_routine): Normalize order of clauses. (cp_finalize_oacc_routine): Call verify_oacc_routine_clauses. diff --git gcc/cp/parser.c gcc/cp/parser.c index 8558a6f..902197b 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -36808,7 +36808,9 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) } /* Process the bind clause, if present. */ - for (tree c = parser->oacc_routine->clauses; c; c = OMP_CLAUSE_CHAIN (c)) + for (tree c = parser->oacc_routine->clauses; + c; + c = OMP_CLAUSE_CHAIN (c)) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_BIND) continue; @@ -36836,34 +36838,39 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) break; } - verify_oacc_routine_clauses (&parser->oacc_routine->clauses, - parser->oacc_routine->loc); - - if (get_oacc_fn_attrib (fndecl)) + int compatible + = verify_oacc_routine_clauses (fndecl, &parser->oacc_routine->clauses, + parser->oacc_routine->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (parser->oacc_routine->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); parser->oacc_routine = NULL; return; } - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + if (compatible > 0) { - error_at (parser->oacc_routine->loc, - "%<#pragma acc routine%> must be applied before %s", - TREE_USED (fndecl) ? "use" : "definition"); - parser->oacc_routine = NULL; - return; } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (parser->oacc_routine->loc, + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + parser->oacc_routine = NULL; + return; + } + + /* Set the routine's level of parallelism. */ + tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses); + replace_oacc_fn_attrib (fndecl, dims); - /* Process the routine's dimension clauses. */ - tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses); - replace_oacc_fn_attrib (fndecl, dims); - - /* Add an "omp declare target" attribute. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl)); + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + parser->oacc_routine->clauses, + DECL_ATTRIBUTES (fndecl)); + } /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git gcc/omp-low.c gcc/omp-low.c index 6604897..b314523 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -12701,13 +12701,18 @@ set_oacc_fn_attrib (tree fn, tree clauses, bool is_kernel, vec *args) /* Verify OpenACC routine clauses. - The chain of clauses returned will contain exactly one clause specifying the - level of parallelism. */ + Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has + already been marked in compatible way, and -1 if incompatible. Upon + returning, the chain of clauses will contain exactly one clause specifying + the level of parallelism. */ -void -verify_oacc_routine_clauses (tree *clauses, location_t loc) +int +verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc, + const char *routine_str) { tree c_level = NULL_TREE; + tree c_bind = NULL_TREE; + tree c_nohost = 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)) @@ -12740,6 +12745,19 @@ verify_oacc_routine_clauses (tree *clauses, location_t loc) c = c_p; } break; + case OMP_CLAUSE_BIND: + /* Don't bother with duplicate clauses at this point. */ + c_bind = c; + break; + case OMP_CLAUSE_NOHOST: + /* Don't bother with duplicate clauses at this point. */ + c_nohost = c; + break; + case OMP_CLAUSE_DEVICE_TYPE: + /* TODO */ + sorry ("% clause not yet supported with %qs", + routine_str); + break; default: gcc_unreachable (); } @@ -12751,6 +12769,126 @@ verify_oacc_routine_clauses (tree *clauses, location_t loc) OMP_CLAUSE_CHAIN (c_level) = *clauses; *clauses = c_level; } + /* In *clauses, we now have exactly one clause specifying the level of + parallelism. */ + + /* Still got some work to do for Fortran... */ + if (fndecl == NULL_TREE) + return 0; + + tree attr + = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)); + if (attr != NULL_TREE) + { + /* If a "#pragma acc routine" has already been applied, just verify + this one for compatibility. */ + /* Collect previous directive's clauses. */ + tree c_level_p = NULL_TREE; + tree c_bind_p = NULL_TREE; + tree c_nohost_p = NULL_TREE; + for (tree c = TREE_VALUE (attr); 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: + gcc_checking_assert (c_level_p == NULL_TREE); + c_level_p = c; + break; + case OMP_CLAUSE_BIND: + /* Don't bother with duplicate clauses at this point. */ + c_bind_p = c; + break; + case OMP_CLAUSE_NOHOST: + /* Don't bother with duplicate clauses at this point. */ + c_nohost_p = c; + break; + case OMP_CLAUSE_DEVICE_TYPE: + /* TODO */ + break; + default: + gcc_unreachable (); + } + gcc_checking_assert (c_level_p != NULL_TREE); + /* ..., and compare to current directive's, which we've already collected + above. */ + tree c_diag; + tree c_diag_p; + /* Matching level of parallelism? */ + if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p)) + { + c_diag = c_level; + c_diag_p = c_level_p; + goto incompatible; + } + /* Matching bind clauses? */ + if ((c_bind == NULL_TREE) != (c_bind_p == NULL_TREE)) + { + c_diag = c_bind; + c_diag_p = c_bind_p; + goto incompatible; + } + /* Matching bind clauses' names? */ + if ((c_bind != NULL_TREE) && (c_bind_p != NULL_TREE)) + { + tree c_bind_name = OMP_CLAUSE_BIND_NAME (c_bind); + tree c_bind_name_p = OMP_CLAUSE_BIND_NAME (c_bind_p); + /* TODO: will/should actually be the trees/strings/string pointers be + identical? */ + if (strcmp (TREE_STRING_POINTER (c_bind_name), + TREE_STRING_POINTER (c_bind_name_p)) != 0) + { + c_diag = c_bind; + c_diag_p = c_bind_p; + goto incompatible; + } + } + /* Matching nohost clauses? */ + if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE)) + { + c_diag = c_nohost; + c_diag_p = c_nohost_p; + goto incompatible; + } + /* Compatible. */ + return 1; + + incompatible: + if (c_diag != NULL_TREE) + error_at (OMP_CLAUSE_LOCATION (c_diag), + "incompatible %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)], + routine_str, fndecl); + else if (c_diag_p != NULL_TREE) + error_at (loc, + "missing %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)], + routine_str, fndecl); + else + gcc_unreachable (); + if (c_diag_p != NULL_TREE) + inform (OMP_CLAUSE_LOCATION (c_diag_p), + "... with %qs clause here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]); + else + { + /* In the front ends, we don't preserve location information for the + OpenACC routine directive itself. However, that of c_level_p + should be close. */ + location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p); + inform (loc_routine, "... without %qs clause near to here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]); + } + /* Incompatible. */ + return -1; + } + + return 0; } /* Process the OpenACC routine's clauses to generate an attribute diff --git gcc/omp-low.h gcc/omp-low.h index c7b7dcb..2602a12 100644 --- gcc/omp-low.h +++ gcc/omp-low.h @@ -31,7 +31,7 @@ extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *); extern void omp_finish_file (void); extern tree omp_member_access_dummy_var (tree); extern void replace_oacc_fn_attrib (tree, tree); -extern void verify_oacc_routine_clauses (tree *, location_t); +extern int verify_oacc_routine_clauses (tree, tree *, location_t, const char *); extern tree build_oacc_routine_dims (tree); extern tree get_oacc_fn_attrib (tree); extern void set_oacc_fn_attrib (tree, tree, bool, vec *); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 0870821..1e12d8f 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,16 @@ 2016-08-17 Thomas Schwinge + * c-c++-common/goacc/routine-1.c: Update. + * c-c++-common/goacc/routine-2.c: Likewise. + * c-c++-common/goacc/routine-nohost-1.c: Likewise. + * g++.dg/goacc/routine-2.C: Likewise. + * c-c++-common/goacc/routine-nohost-2.c: New file. + + * c-c++-common/goacc/oaccdevlow-routine.c: Update. + * c-c++-common/goacc/routine-5.c: Likewise. + * c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise. + * c-c++-common/goacc/routine-level-of-parallelism-2.c: New file. + * c-c++-common/goacc/routine-2.c: Update, and move some test into... * c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c index 7d57921..0b56661 100644 --- gcc/testsuite/c-c++-common/goacc/routine-1.c +++ gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -20,32 +20,75 @@ void seq (void) { } + #pragma acc routine void bind_f_1 (void) { } -#pragma acc routine bind (bind_f_1) +#pragma acc routine +extern void bind_f_1 (void); + +#pragma acc routine (bind_f_1) + + +#pragma acc routine \ + bind (bind_f_1) void bind_f_1_1 (void) { } +#pragma acc routine \ + bind (bind_f_1) +extern void bind_f_1_1 (void); + +#pragma acc routine (bind_f_1_1) \ + bind (bind_f_1) + + /* Non-sensical bind clause, but permitted. */ -#pragma acc routine bind ("bind_f_2") +#pragma acc routine \ + bind ("bind_f_2") void bind_f_2 (void) { } -#pragma acc routine bind ("bind_f_2") +#pragma acc routine \ + bind ("bind_f_2") +extern void bind_f_2 (void); + +#pragma acc routine (bind_f_2) \ + bind ("bind_f_2") + + +#pragma acc routine \ + bind ("bind_f_2") void bind_f_2_1 (void) { } -#pragma acc routine nohost +#pragma acc routine \ + bind ("bind_f_2") +extern void bind_f_2_1 (void); + +#pragma acc routine (bind_f_2_1) \ + bind ("bind_f_2") + + +#pragma acc routine \ + nohost void nohost (void) { } +#pragma acc routine \ + nohost +extern void nohost (void); + +#pragma acc routine (nohost) \ + nohost + + int main () { #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) diff --git gcc/testsuite/c-c++-common/goacc/routine-2.c gcc/testsuite/c-c++-common/goacc/routine-2.c index 7582c86..debf6d7 100644 --- gcc/testsuite/c-c++-common/goacc/routine-2.c +++ gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -26,3 +26,146 @@ extern void bind_3 (void); #pragma acc routine nohost nohost /* { dg-error "too many .nohost. clauses" } */ extern void nohost (void); + + +/* bind clause on first OpenACC routine directive but not on following. */ + +extern void a_bind_f_1 (void); +#pragma acc routine (a_bind_f_1) + + +#pragma acc routine \ + bind (a_bind_f_1) +void a_bind_f_1_1 (void) +{ +} + +#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void a_bind_f_1_1 (void); + +#pragma acc routine (a_bind_f_1_1) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Non-sensical bind clause, but permitted. */ +#pragma acc routine \ + bind ("a_bind_f_2") +void a_bind_f_2 (void) +{ +} + +#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void a_bind_f_2 (void); + +#pragma acc routine (a_bind_f_2) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +#pragma acc routine \ + bind ("a_bind_f_2") +void a_bind_f_2_1 (void) +{ +} + +#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void a_bind_f_2_1 (void); + +#pragma acc routine (a_bind_f_2_1) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* No bind clause on first OpenACC routine directive, but on following. */ + +#pragma acc routine +extern void b_bind_f_1 (void); + + +#pragma acc routine +void b_bind_f_1_1 (void) +{ +} + +#pragma acc routine \ + bind (b_bind_f_1) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void b_bind_f_1_1 (void); + +#pragma acc routine (b_bind_f_1_1) \ + bind (b_bind_f_1) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Non-sensical bind clause, but permitted. */ +#pragma acc routine +void b_bind_f_2 (void) +{ +} + +#pragma acc routine \ + bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void b_bind_f_2 (void); + +#pragma acc routine (b_bind_f_2) \ + bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +#pragma acc routine +void b_bind_f_2_1 (void) +{ +} + +#pragma acc routine \ + bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void b_bind_f_2_1 (void); + +#pragma acc routine (b_bind_f_2_1) \ + bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Non-matching bind clauses. */ + +#pragma acc routine +void c_bind_f_1a (void) +{ +} + +#pragma acc routine +extern void c_bind_f_1b (void); + + +#pragma acc routine \ + bind (c_bind_f_1a) +void c_bind_f_1_1 (void) +{ +} + +#pragma acc routine \ + bind (c_bind_f_1b) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void c_bind_f_1_1 (void); + +#pragma acc routine (c_bind_f_1_1) \ + bind (c_bind_f_1b) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Non-sensical bind clause, but permitted. */ +#pragma acc routine \ + bind ("c_bind_f_2") +void c_bind_f_2 (void) +{ +} + +#pragma acc routine \ + bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void c_bind_f_2 (void); + +#pragma acc routine (c_bind_f_2) \ + bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +#pragma acc routine \ + bind ("c_bind_f_2") +void c_bind_f_2_1 (void) +{ +} + +#pragma acc routine \ + bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void c_bind_f_2_1 (void); + +#pragma acc routine (c_bind_f_2_1) \ + bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ diff --git gcc/testsuite/c-c++-common/goacc/routine-5.c gcc/testsuite/c-c++-common/goacc/routine-5.c index 17fe67c..f10651d 100644 --- gcc/testsuite/c-c++-common/goacc/routine-5.c +++ gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -150,61 +150,19 @@ void f_static_assert(); #pragma acc routine __extension__ extern void ex1(); -#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */ +#pragma acc routine (ex1) worker /* { dg-error "has already been marked as an accelerator routine" } */ #pragma acc routine __extension__ __extension__ __extension__ __extension__ __extension__ void ex2() { } -#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */ +#pragma acc routine (ex2) worker /* { dg-error "has already been marked as an accelerator routine" } */ #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ __extension__ int ex3; #pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */ -/* "#pragma acc routine" already applied. */ - -extern void fungsi_1(); -#pragma acc routine(fungsi_1) gang -#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ - -#pragma acc routine seq -extern void fungsi_2(); -#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ - -#pragma acc routine vector -extern void fungsi_3(); -#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */ -void fungsi_3() -{ -} - -extern void fungsi_4(); -#pragma acc routine (fungsi_4) worker -#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */ -void fungsi_4() -{ -} - -#pragma acc routine gang -void fungsi_5() -{ -} -#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */ - -#pragma acc routine seq -void fungsi_6() -{ -} -#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */ -extern void fungsi_6(); - - /* "#pragma acc routine" must be applied before. */ void Bar (); diff --git gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c index fc8bc3c..8f45499 100644 --- gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c +++ gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c @@ -42,10 +42,10 @@ void s_2 (void) void g_3 (void) { } -#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (g_3) \ gang \ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ @@ -53,10 +53,10 @@ 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (w_3) \ worker \ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ @@ -66,10 +66,10 @@ extern void w_3 (void); void v_3 (void) { } -#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (v_3) \ vector \ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ @@ -77,10 +77,10 @@ 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (s_3) \ seq \ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ @@ -91,12 +91,12 @@ extern void s_3 (void); 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -108,12 +108,12 @@ extern void w_4 (void); 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -127,12 +127,12 @@ extern void w_4 (void); void v_4 (void) { } -#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -146,12 +146,12 @@ void v_4 (void) void s_4 (void) { } -#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -169,7 +169,7 @@ void s_4 (void) void g_5 (void) { } -#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \ +#pragma acc routine (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 *-*-* } 174 } */ \ @@ -177,7 +177,7 @@ void g_5 (void) /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 176 } */ \ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 178 } */ -#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \ +#pragma acc routine (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 *-*-* } 182 } */ \ @@ -195,7 +195,7 @@ void g_5 (void) gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 195 } */ extern void w_5 (void); -#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \ +#pragma acc routine (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 *-*-* } 200 } */ \ @@ -203,7 +203,7 @@ extern void w_5 (void); /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 202 } */ \ gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 204 } */ -#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \ +#pragma acc routine (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 *-*-* } 208 } */ \ @@ -221,7 +221,7 @@ extern void w_5 (void); gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 221 } */ extern void v_5 (void); -#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \ +#pragma acc routine (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 *-*-* } 226 } */ \ @@ -229,7 +229,7 @@ extern void v_5 (void); /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 228 } */ \ gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 230 } */ -#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \ +#pragma acc routine (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 *-*-* } 234 } */ \ @@ -247,7 +247,7 @@ extern void s_5 (void); /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 246 } */ \ gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 248 } */ -#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \ +#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 *-*-* } 252 } */ \ @@ -255,7 +255,7 @@ extern void s_5 (void); /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 254 } */ \ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 256 } */ -#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \ +#pragma acc routine (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 *-*-* } 260 } */ \ @@ -263,3 +263,188 @@ extern void s_5 (void); /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 262 } */ \ gang gang /* { dg-error "too many 'gang' clauses" } */ \ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 264 } */ + + +/* Like the *_5 tests, but with the order of clauses changed in the second and + following routine directives for the specific *_5 function. */ + +#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 *-*-* } 273 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 275 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 277 } */ +void g_6 (void) +{ +} +#pragma acc routine (g_6) \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 283 } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 285 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 287 } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 289 } */ +#pragma acc routine (g_6) \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 292 } */ \ + gang gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 294 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 296 } */ \ + worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 298 } */ + +#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 *-*-* } 303 } */ \ + vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 305 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 307 } */ +extern void w_6 (void); +#pragma acc routine (w_6) \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 311 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 313 } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 315 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 317 } */ +#pragma acc routine (w_6) \ + seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 320 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 322 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 324 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 326 } */ + +#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 *-*-* } 331 } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 333 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 335 } */ +extern void v_6 (void); +#pragma acc routine (v_6) \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 339 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 341 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 343 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 345 } */ +#pragma acc routine (v_6) \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 348 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 350 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 352 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 354 } */ + +extern void s_6 (void); +#pragma acc routine (s_6) \ + 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 *-*-* } 360 } */ \ + worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 362 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 364 } */ +#pragma acc routine (s_6) \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 367 } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 369 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 371 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 373 } */ +#pragma acc routine (s_6) \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 376 } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 378 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 380 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 382 } */ + + +/* Like the *_6 tests, but without all the duplicate clauses, so that the + routine directives are valid in isolation. */ + +#pragma acc routine \ + gang +void g_7 (void) +{ +} +#pragma acc routine (g_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (g_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + worker +extern void w_7 (void); +#pragma acc routine (w_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (w_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + vector +extern void v_7 (void); +#pragma acc routine (v_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (v_7) \ + gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +extern void s_7 (void); +#pragma acc routine (s_7) \ + seq +#pragma acc routine (s_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_7) \ + worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Test cases for implicit seq clause. */ + +#pragma acc routine \ + gang +void g_8 (void) +{ +} +#pragma acc routine (g_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + worker +extern void w_8 (void); +#pragma acc routine (w_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + vector +extern void v_8 (void); +#pragma acc routine (v_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +extern void s_8 (void); +#pragma acc routine (s_8) +#pragma acc routine (s_8) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_8) \ + gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_8) \ + worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ diff --git gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c new file mode 100644 index 0000000..1803997 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c @@ -0,0 +1,73 @@ +/* Test various aspects of clauses specifying compatible levels of + parallelism with the OpenACC routine directive. The Fortran counterpart is + ../../gfortran.dg/goacc/routine-level-of-parallelism-2.f. */ + +#pragma acc routine gang +void g_1 (void) /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ +/* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "worker partitioned" { xfail *-*-* } 6 } */ +/* { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "worker partitioned" { xfail *-*-* } 6 } */ +{ +} +#pragma acc routine (g_1) gang +#pragma acc routine (g_1) gang + + +extern void w_1 (void); +#pragma acc routine (w_1) worker +#pragma acc routine (w_1) worker +#pragma acc routine (w_1) worker + + +#pragma acc routine vector +extern void v_1 (void); +#pragma acc routine (v_1) vector +#pragma acc routine (v_1) vector + + +/* Also test the implicit seq clause. */ + +#pragma acc routine seq +extern void s_1_1 (void); +#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) seq +#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) seq + +#pragma acc routine +extern void s_1_2 (void); +#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) seq +#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) seq + +extern void s_2_1 (void); +#pragma acc routine (s_2_1) seq +#pragma acc routine (s_2_1) +#pragma acc routine (s_2_1) seq +#pragma acc routine (s_2_1) +#pragma acc routine (s_2_1) seq + +extern void s_2_2 (void); +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) seq +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) seq + +#pragma acc routine seq +void s_3_1 (void) +{ +} +#pragma acc routine (s_3_1) +#pragma acc routine (s_3_1) seq +#pragma acc routine (s_3_1) +#pragma acc routine (s_3_1) seq + +#pragma acc routine +void s_3_2 (void) +{ +} +#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) seq +#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) seq diff --git gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c index 9baa56c..17d6b03 100644 --- gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c +++ gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -9,13 +9,27 @@ int THREE(void) return 3; } +#pragma acc routine (THREE) nohost + +#pragma acc routine nohost +extern int THREE(void); + + #pragma acc routine nohost extern void NOTHING(void); +#pragma acc routine (NOTHING) nohost + void NOTHING(void) { } +#pragma acc routine nohost +extern void NOTHING(void); + +#pragma acc routine (NOTHING) nohost + + extern float ADD(float, float); #pragma acc routine (ADD) nohost @@ -25,4 +39,10 @@ float ADD(float x, float y) return x + y; } +#pragma acc routine nohost +extern float ADD(float, float); + +#pragma acc routine (ADD) nohost + + /* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccdevlow" } } */ diff --git gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c new file mode 100644 index 0000000..7402bfc --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c @@ -0,0 +1,97 @@ +/* Test invalid usage of the nohost clause for OpenACC routine directive. + Exercising different variants for declaring routines. */ + +#pragma acc routine +int THREE_1(void) +{ + return 3; +} + +#pragma acc routine (THREE_1) \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern int THREE_1(void); + + +#pragma acc routine +extern void NOTHING_1(void); + +#pragma acc routine (NOTHING_1) \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +void NOTHING_1(void) +{ +} + +#pragma acc routine \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void NOTHING_1(void); + +#pragma acc routine (NOTHING_1) \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +extern float ADD_1(float, float); + +#pragma acc routine (ADD_1) + +float ADD_1(float x, float y) +{ + return x + y; +} + +#pragma acc routine \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_1\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */ +extern float ADD_1(float, float); + +#pragma acc routine (ADD_1) \ + nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_1\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */ + + +/* The same again, but with/without nohost reversed. */ + +#pragma acc routine \ + nohost +int THREE_2(void) +{ + return 3; +} + +#pragma acc routine (THREE_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern int THREE_2(void); + + +#pragma acc routine \ + nohost +extern void NOTHING_2(void); + +#pragma acc routine (NOTHING_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +void NOTHING_2(void) +{ +} + +#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +extern void NOTHING_2(void); + +#pragma acc routine (NOTHING_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +extern float ADD_2(float, float); + +#pragma acc routine (ADD_2) \ + nohost + +float ADD_2(float x, float y) +{ + return x + y; +} + +#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_2\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */ +extern float ADD_2(float, float); + +#pragma acc routine (ADD_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_2\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */ diff --git gcc/testsuite/g++.dg/goacc/routine-2.C gcc/testsuite/g++.dg/goacc/routine-2.C index cbee487..939dfba 100644 --- gcc/testsuite/g++.dg/goacc/routine-2.C +++ gcc/testsuite/g++.dg/goacc/routine-2.C @@ -58,12 +58,18 @@ broken_mult2 (float a, float b) { return a + b; } +#pragma acc routine (broken_mult2) bind("mult") +#pragma acc routine bind("mult") +extern float broken_mult2 (float, float); #pragma acc routine bind(sum2) /* { dg-error "'sum2' has not been declared" } */ int broken_mult3 (int a, int b); #pragma acc routine bind(foo::sub) int broken_mult4 (int a, int b); +#pragma acc routine (broken_mult4) bind(foo::sub) +#pragma acc routine bind(foo::sub) +extern int broken_mult4 (int a, int b); namespace bar { @@ -72,6 +78,9 @@ namespace bar { return a * b; } +#pragma acc routine (working_mult) bind (mult) +#pragma acc routine bind (mult) + float working_mult (float, float); } #pragma acc routine