From patchwork Fri Nov 11 23:44:29 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 81908 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp9897qge; Fri, 11 Nov 2016 15:45:11 -0800 (PST) X-Received: by 10.99.169.25 with SMTP id u25mr8823958pge.6.1478907911199; Fri, 11 Nov 2016 15:45:11 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id 6si12615251pfl.234.2016.11.11.15.45.10 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 11 Nov 2016 15:45:11 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-441212-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-441212-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-441212-patch=linaro.org@gcc.gnu.org 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:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=muAU+TbmQVfZfhLGVV4eOL92DXSX7GZ4ldzoHQw6eGNCVfpfVf mXupk5uKEaevwa0iuJTzHzINYaxgEItJZAo3vzLAFXt5CFcmNa8h04YnjbyihqhY 61GITBPBPm5Er8RtrncC9bcDGZvVmaDMP7UmiTz7ogzRa76XYNRPB82Yc= 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:cc:message-id:date:mime-version:content-type; s= default; bh=GsBQuuFg6KUJKT7sZEgDLp04s1M=; b=VqOKbL7RnnHo1Ke7qmy8 Bx2DdGee2RG6f6rKefTdzMHwmB1LMRbRUOrWwiVgXbeTPu1xgreDPr1O+lwDz8TR 9P3I5voDI/KLYHhT6CFK10HM9HM0RMQT8NKqh1+oW9oIJCL0KenVJUy6Qp57x9dD aOdUJ3C54AN/DtCndRgUWCg= Received: (qmail 57683 invoked by alias); 11 Nov 2016 23:44:49 -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 57655 invoked by uid 89); 11 Nov 2016 23:44:47 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=NOTHING, TWO, 2366, 236, 6 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; Fri, 11 Nov 2016 23:44:36 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c5LUo-0002wL-Dw from Cesar_Philippidis@mentor.com ; Fri, 11 Nov 2016 15:44:34 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Fri, 11 Nov 2016 15:44:31 -0800 From: Cesar Philippidis Subject: [PATCH] OpenACC routines -- test cases To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek CC: Thomas Schwinge Message-ID: <6b840a02-5355-5e87-98e8-21968429589c@codesourcery.com> Date: Fri, 11 Nov 2016 15:44:29 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) This patch contains new and adjusted, runtime and compiler test cases for the new OpenACC routine functionality. Is this ok for trunk? Cesar 2016-11-11 Cesar Philippidis Thomas Schwinge gcc/testsuite/ * c-c++-common/goacc/routine-1.c: Add more coverage. * c-c++-common/goacc/routine-2.c: Likewise. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-1.c: New test. * c-c++-common/goacc/routine-level-of-parallelism-2.c: New test. * c-c++-common/goacc/routine-nohost-1.c: New test. * c-c++-common/goacc/routine-nohost-2.c: New test. * gfortran.dg/goacc/fixed-1.f: Add more coverage. * gfortran.dg/goacc/loop-5.f95: Likewise. * gfortran.dg/goacc/routine-6.f90: Likewise. * gfortran.dg/goacc/routine-7.f90: New test. * gfortran.dg/goacc/routine-9.f90: New test. * gfortran.dg/goacc/routine-nested-parallelism.f: New test. * gfortran.dg/goacc/routine-nested-parallelism.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/routine-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/routine-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New test. * testsuite/libgomp.oacc-fortran/abort-1.f90: Add more coverage. * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. * testsuite/libgomp.oacc-fortran/routine-6.f90: New test. * testsuite/libgomp.oacc-fortran/routine-7.f90: Add more coverage. * testsuite/libgomp.oacc-fortran/routine-8.f90: New test. * testsuite/libgomp.oacc-fortran/vector-routine.f90: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index a5e0d69..43434a3 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -1,3 +1,4 @@ +/* Test valid use of clauses with routine. */ #pragma acc routine gang void gang (void) @@ -19,15 +20,88 @@ void seq (void) { } -int main () + +#pragma acc routine +void bind_f_1 (void) +{ +} + +#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") +void bind_f_2 (void) +{ +} + +#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 \ + 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) { gang (); worker (); vector (); seq (); + bind_f_1 (); + bind_f_1_1 (); + bind_f_2 (); + bind_f_2_1 (); + nohost (); } return 0; diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c index fc5eb11..debf6d7 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -1,21 +1,171 @@ -#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */ -void gang (void) +/* Test invalid use of clauses with OpenACC routine. */ + +extern float F; +#pragma acc routine bind (F) /* { dg-error ".F. does not refer to a function" } */ +extern void F_1 (void); + +typedef int T; +#pragma acc routine bind (T) /* { dg-error ".T. does not refer to a function" } */ +extern void T_1 (void); + +#pragma acc routine (nothing) gang /* { dg-error ".nothing. has not been declared" } */ + +#pragma acc routine bind (bind_0) /* { dg-error ".bind_0. has not been declared" }*/ +extern void bind_0 (void); + +extern void a(void), b(void); + +#pragma acc routine bind(a) bind(b) /* { dg-error "too many .bind. clauses" } */ +extern void bind_1 (void); + +#pragma acc routine bind(a) bind("b") /* { dg-error "too many .bind. clauses" } */ +extern void bind_2 (void); + +#pragma acc routine bind("a") bind("b") /* { dg-error "too many .bind. clauses" } */ +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 worker vector /* { dg-error "multiple loop axes" } */ -void worker (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 vector seq /* { dg-error "multiple loop axes" } */ -void vector (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 seq gang /* { dg-error "multiple loop axes" } */ -void seq (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 (nothing) gang /* { dg-error "not been declared" } */ +#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 a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index 17fe67c..f10651d 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/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 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 0000000..8f45499 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c @@ -0,0 +1,450 @@ +/* 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) \ + gang \ + seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ +#pragma acc routine (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) \ + worker \ + gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ +#pragma acc routine (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) \ + vector \ + worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ +#pragma acc routine (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) \ + seq \ + vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ +#pragma acc routine (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) \ + 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) \ + 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) \ + 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) \ + 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) \ + 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) \ + 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) \ + 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) \ + 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 *-*-* } 163 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 165 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 167 } */ +void g_5 (void) +{ +} +#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 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { 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) \ + 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 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 184 } */ \ + worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 186 } */ + +#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 *-*-* } 191 } */ \ + vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 193 } */ \ + 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) \ + 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 } */ \ + seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { 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) \ + 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 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 210 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 212 } */ + +#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 *-*-* } 217 } */ \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 219 } */ \ + 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) \ + 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 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { 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) \ + 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 } */ \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 236 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 238 } */ + +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 *-*-* } 244 } */ \ + worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { 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) \ + 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 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { 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) \ + 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 } */ \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { 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 a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c new file mode 100644 index 0000000..1217397 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c @@ -0,0 +1,71 @@ +/* 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) +{ +} +#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 a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c new file mode 100644 index 0000000..17d6b03 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -0,0 +1,48 @@ +/* Test the nohost clause for OpenACC routine directive. Exercising different + variants for declaring routines. */ + +/* { dg-additional-options "-fdump-tree-oaccdevlow" } */ + +#pragma acc routine nohost +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 + +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 a/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c new file mode 100644 index 0000000..7402bfc --- /dev/null +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f index 6a454190..0c0fb98 100644 --- a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f @@ -1,3 +1,5 @@ +!$ACC ROUTINE(ABORT) SEQ + INTEGER :: ARGC ARGC = COMMAND_ARGUMENT_COUNT () diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 index d059cf7..fe137d5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 @@ -93,9 +93,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc loop vector tile(*) DO i = 1,10 ENDDO @@ -129,9 +126,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc loop vector tile(*) DO i = 1,10 ENDDO @@ -242,9 +236,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc kernels loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc kernels loop vector tile(*) DO i = 1,10 ENDDO @@ -333,9 +324,6 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc parallel loop tile(-1) ! { dg-warning "must be positive" } - do i = 1,10 - enddo !$acc parallel loop vector tile(*) DO i = 1,10 ENDDO diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 index 10943cf..ee43935 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 @@ -5,7 +5,6 @@ contains subroutine subr5 (x) implicit none !$acc routine (subr5) - !$acc routine (m1int) ! { dg-error "invalid function name" } integer, intent(inout) :: x if (x < 1) then x = 1 @@ -26,7 +25,6 @@ program main end interface integer, parameter :: n = 10 integer :: a(n), i - !$acc routine (subr1) ! { dg-error "invalid function name" } external :: subr2 !$acc routine (subr2) @@ -56,7 +54,6 @@ subroutine subr1 (x) end subroutine subr1 subroutine subr2 (x) - !$acc routine (subr1) ! { dg-error "invalid function name" } integer, intent(inout) :: x if (x < 1) then x = 1 @@ -86,7 +83,6 @@ subroutine subr4 (x) end subroutine subr4 subroutine subr10 (x) - !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" } integer, intent(inout) :: x if (x < 1) then x = 1 diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-7.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90 new file mode 100644 index 0000000..cb45079 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90 @@ -0,0 +1,70 @@ +! Test acc routines inside modules. + +! { dg-additional-options "-O0" } + +module routines + integer a +contains + subroutine vector + implicit none + !$acc routine vector + end subroutine vector + + subroutine worker + implicit none + !$acc routine worker + end subroutine worker + + subroutine gang + implicit none + !$acc routine gang + end subroutine gang + + subroutine seq + implicit none + !$acc routine seq + end subroutine seq +end module routines + +program main + use routines + implicit none + + integer :: i + + !$acc parallel loop gang + do i = 1, 10 + call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call worker + call vector + call seq + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, 10 + call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call worker ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call vector + call seq + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, 10 + call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call worker ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call vector ! { dg-error "routine call uses same OpenACC parallelism as containing loop" } + call seq + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, 10 + call gang + call worker + call vector + call seq + end do + !$acc end parallel loop +end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 new file mode 100644 index 0000000..590e594 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 @@ -0,0 +1,96 @@ +! Check for late resolver errors caused by invalid ACC ROUTINE +! directives. + +module m + integer m1int +contains + subroutine subr5 (x) + implicit none + integer extfunc + !$acc routine (subr5) + !$acc routine (extfunc) + !$acc routine (m1int) ! { dg-error "invalid function name" } + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + x = x * x - 1 + extfunc(2) + end if + end subroutine subr5 +end module m + +program main + implicit none + interface + function subr6 (x) + integer, intent (in) :: x + integer :: subr6 + end function subr6 + end interface + integer, parameter :: n = 10 + integer :: a(n), i + !$acc routine (subr1) ! { dg-error "invalid function name" } + external :: subr2 + !$acc routine (subr2) + + external :: R1, R2 + !$acc routine (R1) + !$acc routine (R2) + + !$acc parallel + !$acc loop + do i = 1, n + call subr1 (i) + call subr2 (i) + end do + !$acc end parallel +end program main + +subroutine subr1 (x) + !$acc routine + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + x = x * x - 1 + end if +end subroutine subr1 + +subroutine subr2 (x) + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + x = x * x - 1 + end if +end subroutine subr2 + +subroutine subr3 (x) + !$acc routine (subr3) + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + call subr4 (x) + end if +end subroutine subr3 + +subroutine subr4 (x) + !$acc routine (subr4) + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + x = x * x - 1 + end if +end subroutine subr4 + +subroutine subr10 (x) + !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" } + integer, intent(inout) :: x + if (x < 1) then + x = 1 + else + x = x * x - 1 + end if +end subroutine subr10 diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f new file mode 100644 index 0000000..d1304c6 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f @@ -0,0 +1,340 @@ +! Validate calls to ACC ROUTINES. Ensure that the loop containing the +! call has sufficient parallelism to for the routine. + + subroutine sub + implicit none + integer, parameter :: n = 100 + integer :: a(n), i, j + external gangr, workerr, vectorr, seqr +!$acc routine (gangr) gang +!$acc routine (workerr) worker +!$acc routine (vectorr) vector +!$acc routine (seqr) seq + +! +! Test subroutine calls inside nested loops. +! + +!$acc parallel loop + do i = 1, n + !$acc loop + do j = 1, n + call workerr (a, n) + end do + end do +!$acc end parallel loop + +!$acc parallel loop + do i = 1, n +!$acc loop gang + do j = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + end do +!$acc end parallel loop + +! +! Test calls to seq routines +! + +!$acc parallel loop + do i = 1, n + call seqr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + call seqr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + call seqr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + call seqr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + call seqr (a, n) + end do +!$acc end parallel loop + +! +! Test calls to gang routines +! + +!$acc parallel loop + do i = 1, n + call gangr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +! +! Test calls to worker routines +! + +!$acc parallel loop + do i = 1, n + call workerr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + call workerr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +! +! Test calls to vector routines +! + +!$acc parallel loop + do i = 1, n + call vectorr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + call vectorr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + call vectorr (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + end subroutine sub + + subroutine func + implicit none + integer, parameter :: n = 100 + integer :: a(n), i, j + integer gangf, workerf, vectorf, seqf +!$acc routine (gangf) gang +!$acc routine (workerf) worker +!$acc routine (vectorf) vector +!$acc routine (seqf) seq + +! +! Test subroutine calls inside nested loops. +! + +!$acc parallel loop + do i = 1, n +!$acc loop + do j = 1, n + a(1) = workerf (a, n) + end do + end do +!$acc end parallel loop + +!$acc parallel loop + do i = 1, n +!$acc loop gang + do j = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + end do +!$acc end parallel loop + +! +! Test calls to seq routines +! + +!$acc parallel loop + do i = 1, n + a(1) = seqf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + a(1) = seqf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + a(1) = seqf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + a(1) = seqf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + a(1) = seqf (a, n) + end do +!$acc end parallel loop + +! +! Test calls to gang routines +! + +!$acc parallel loop + do i = 1, n + a(1) = gangf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +! +! Test calls to worker routines +! + +!$acc parallel loop + do i = 1, n + a(1) = workerf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + a(1) = workerf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +! +! Test calls to vector routines +! + +!$acc parallel loop + do i = 1, n + a(1) = vectorf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop gang + do i = 1, n + a(1) = vectorf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop worker + do i = 1, n + a(1) = vectorf (a, n) + end do +!$acc end parallel loop + +!$acc parallel loop vector + do i = 1, n + a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + +!$acc parallel loop seq + do i = 1, n + a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do +!$acc end parallel loop + end subroutine func diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90 new file mode 100644 index 0000000..94e0464 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90 @@ -0,0 +1,340 @@ +! Validate calls to ACC ROUTINES. Ensure that the loop containing the +! call has sufficient parallelism to for the routine. + +subroutine sub + implicit none + integer, parameter :: n = 100 + integer :: a(n), i, j + external gangr, workerr, vectorr, seqr + !$acc routine (gangr) gang + !$acc routine (workerr) worker + !$acc routine (vectorr) vector + !$acc routine (seqr) seq + + ! + ! Test subroutine calls inside nested loops. + ! + + !$acc parallel loop + do i = 1, n + !$acc loop + do j = 1, n + call workerr (a, n) + end do + end do + !$acc end parallel loop + + !$acc parallel loop + do i = 1, n + !$acc loop gang + do j = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + end do + !$acc end parallel loop + + ! + ! Test calls to seq routines + ! + + !$acc parallel loop + do i = 1, n + call seqr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + call seqr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + call seqr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + call seqr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + call seqr (a, n) + end do + !$acc end parallel loop + + ! + ! Test calls to gang routines + ! + + !$acc parallel loop + do i = 1, n + call gangr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + ! + ! Test calls to worker routines + ! + + !$acc parallel loop + do i = 1, n + call workerr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + call workerr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + ! + ! Test calls to vector routines + ! + + !$acc parallel loop + do i = 1, n + call vectorr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + call vectorr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + call vectorr (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop +end subroutine sub + +subroutine func + implicit none + integer, parameter :: n = 100 + integer :: a(n), i, j + integer gangf, workerf, vectorf, seqf + !$acc routine (gangf) gang + !$acc routine (workerf) worker + !$acc routine (vectorf) vector + !$acc routine (seqf) seq + + ! + ! Test subroutine calls inside nested loops. + ! + + !$acc parallel loop + do i = 1, n + !$acc loop + do j = 1, n + a(1) = workerf (a, n) + end do + end do + !$acc end parallel loop + + !$acc parallel loop + do i = 1, n + !$acc loop gang + do j = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + end do + !$acc end parallel loop + + ! + ! Test calls to seq routines + ! + + !$acc parallel loop + do i = 1, n + a(1) = seqf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + a(1) = seqf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + a(1) = seqf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + a(1) = seqf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + a(1) = seqf (a, n) + end do + !$acc end parallel loop + + ! + ! Test calls to gang routines + ! + + !$acc parallel loop + do i = 1, n + a(1) = gangf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + ! + ! Test calls to worker routines + ! + + !$acc parallel loop + do i = 1, n + a(1) = workerf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + a(1) = workerf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + ! + ! Test calls to vector routines + ! + + !$acc parallel loop + do i = 1, n + a(1) = vectorf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, n + a(1) = vectorf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, n + a(1) = vectorf (a, n) + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, n + a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, n + a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" } + end do + !$acc end parallel loop +end subroutine func diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c new file mode 100644 index 0000000..2cdd6bf --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c @@ -0,0 +1,33 @@ +/* At -O0, we do get the expected "undefined reference to `foo'" link-time + error message (but the check needs to be done differently; compare to + routine-nohost-1.c), but for -O2 we don't; presumably because the function + gets inlined. + { dg-xfail-if "TODO" { *-*-* } { "-O0" } { "" } } */ + +#include + +#pragma acc routine nohost +int +foo (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * n; +} + +int +main() +{ + int a, n = 10; + +#pragma acc parallel copy (a, n) + { + a = foo (n); + } + + if (a != n * n) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c new file mode 100644 index 0000000..4e34f78 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c @@ -0,0 +1,64 @@ + +/* { dg-do run } */ +/* { dg-warning "TODO" "implicit" { xfail *-*-* } 17 } */ +/* { dg-warning "TODO" "implicit" { xfail *-*-* } 27 } */ +/* { dg-xfail-if "unresolved symbol" { *-*-* } } */ + +#include + +#pragma acc routine bind (foo) +int +subr1 (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * foo (n - 1); +} + +#pragma acc routine bind ("bar") +int +subr2 (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * bar (n - 1); +} + +int +main() +{ + int *a, i, n = 10; + + a = (int *)malloc (sizeof (int) * n); + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = foo (i); + } + + for (i = 0; i < n; i++) + if (a[i] != subr1 (i)) + abort (); + + for (i = 0; i < n; i++) + a[i] = 0; + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = bar (i); + } + + for (i = 0; i < n; i++) + if (a[i] != subr2 (i)) + abort (); + + free (a); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c new file mode 100644 index 0000000..b991bb1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c @@ -0,0 +1,105 @@ +/* Test the bind and nohost clauses for OpenACC routine directive. */ + +/* TODO. Function inlining and the OpenACC bind clause do not yet get on well + with one another. + { dg-additional-options "-fno-inline" } */ + +/* TODO. C works, but for C++ we get: "lto1: internal compiler error: in + ipa_propagate_frequency". + { dg-xfail-if "TODO" { *-*-* } } */ + +#include + +/* "MINUS_TWO" is the device variant for function "TWO". Similar for "THREE", + and "FOUR". Exercising different variants for declaring routines. */ + +#pragma acc routine nohost +extern int MINUS_TWO(void); + +int MINUS_TWO(void) +{ + if (!acc_on_device(acc_device_not_host)) + __builtin_abort(); + return -2; +} + +extern int TWO(void); +#pragma acc routine (TWO) bind(MINUS_TWO) + +int TWO(void) +{ + if (acc_on_device(acc_device_not_host)) + __builtin_abort(); + return 2; +} + + +#pragma acc routine nohost +int MINUS_THREE(void) +{ + if (!acc_on_device(acc_device_not_host)) + __builtin_abort(); + return -3; +} + +#pragma acc routine bind(MINUS_THREE) +extern int THREE(void); + +int THREE(void) +{ + if (acc_on_device(acc_device_not_host)) + __builtin_abort(); + return 3; +} + + +/* Due to using a string in the bind clause, we don't need "MINUS_FOUR" in + scope here. */ +#pragma acc routine bind("MINUS_FOUR") +int FOUR(void) +{ + if (acc_on_device(acc_device_not_host)) + __builtin_abort(); + return 4; +} + +extern int MINUS_FOUR(void); +#pragma acc routine (MINUS_FOUR) nohost + +int MINUS_FOUR(void) +{ + if (!acc_on_device(acc_device_not_host)) + __builtin_abort(); + return -4; +} + + +int main() +{ + int x2, x3, x4; + +#pragma acc parallel copyout(x2, x3, x4) if(0) + { + x2 = TWO(); + x3 = THREE(); + x4 = FOUR(); + } + if (x2 != 2 || x3 != 3 || x4 != 4) + __builtin_abort(); + +#pragma acc parallel copyout(x2, x3, x4) + { + x2 = TWO(); + x3 = THREE(); + x4 = FOUR(); + } +#ifdef ACC_DEVICE_TYPE_host + if (x2 != 2 || x3 != 3 || x4 != 4) + __builtin_abort(); +#else + if (x2 != -2 || x3 != -3 || x4 != -4) + __builtin_abort(); +#endif + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c new file mode 100644 index 0000000..365af93 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c @@ -0,0 +1,18 @@ +/* { dg-do link } */ + +extern int three (void); + +#pragma acc routine (three) nohost +__attribute__((noinline)) +int three(void) +{ + return 3; +} + +int main(void) +{ + return (three() == 3) ? 0 : 1; +} + +/* Expecting link to fail; "undefined reference to `three'" (or similar). + { dg-excess-errors "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 index b38303d..48ebc38 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 @@ -1,5 +1,6 @@ program main implicit none + !$acc routine(abort) seq print *, "CheCKpOInT" !$acc parallel diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f index a19045b..cbd1dd9 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f @@ -6,6 +6,7 @@ USE OPENACC IMPLICIT NONE +!$ACC ROUTINE(ABORT) SEQ !Host. diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 new file mode 100644 index 0000000..1bae09c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 @@ -0,0 +1,28 @@ +! { dg-do run } +! { dg-xfail-if "TODO" { *-*-* } } + +program main + integer :: a, n + + n = 10 + + !$acc parallel copy (a, n) + a = foo (n) + !$acc end parallel + + if (a .ne. n * n) call abort + +contains + +function foo (n) result (rc) + !$acc routine nohost + + integer, intent (in) :: n + integer :: rc + + rc = n * n + +end function + +end program main + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 index 200188e..ef2ff04 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 @@ -100,7 +100,7 @@ subroutine gang (a) integer, intent (inout) :: a(N) integer :: i - !$acc loop gang + !$acc loop gang worker vector do i = 1, N a(i) = a(i) - i end do diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 new file mode 100644 index 0000000..5c58b43 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 @@ -0,0 +1,62 @@ + +! { dg-do run } +! { dg-error "Invalid" "TODO" { xfail *-*-* } 51 } + +program main + integer, parameter :: n = 10 + integer :: a(n) + integer :: i + + !$acc parallel copy (a) vector_length (5) + !$acc loop + do i = 1, n + a(i) = foo (i); + end do + !$acc end parallel + + do i = 1, n + if (a(i) .ne. subr1 (i)) call abort + end do + + do i = 1, n + a(i) = 0 + end do + + !$acc parallel copy (a) vector_length (5) + !$acc loop + do i = 1, n + a(i) = bar (i); + end do + !$acc end parallel + + do i = 1, n + if (a(i) .ne. subr2 (i)) call abort + end do + +contains + +function subr1 (n) result (rc) + !$acc routine bind (foo) + integer :: n, rc + + if ((n .eq. 0) .or. (n .eq. 1)) then + rc = 1 + else + rc = n * foo (n - 1); + end if + +end function + +function subr2 (n) result (rc) + !$acc routine bind ("bar") + integer :: n, rc + + if ((n .eq. 0) .or. (n .eq. 1)) then + rc = 1 + else + rc = n * bar (n - 1); + end if + +end function + +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90 b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90 new file mode 100644 index 0000000..1edcee4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90 @@ -0,0 +1,41 @@ +! { dg-do run } + +module param + integer, parameter :: N = 32 +end module param + +program main + use param + integer :: i + integer :: a(N) + + do i = 1, N + a(i) = i + end do + + !$acc parallel copy (a) + !$acc loop worker + do i = 1, N + call vector (a) + end do + !$acc end parallel + + do i = 1, N + if (a(i) .ne. 0) call abort + end do + +contains + + subroutine vector (a) + !$acc routine vector + integer, intent (inout) :: a(N) + integer :: i + + !$acc loop vector + do i = 1, N + a(i) = a(i) - a(i) + end do + +end subroutine vector + +end program main