The language regarding OpenACC routines have changed in OpenACC 2.5 such that the end user must explicitly specify one of gang, worker, vector or seq partitioning. I guess some other compiler need those directives to generate specialized versions of those functions accordingly. However, GCC currently falls back to 'seq' partitioning, and that seems to be a reasonable compromise in most cases.
This patch teaches the FE's how to emit a warning whenever a routine directive doesn't include a partitioning clause. The rationale for making this a warning vs an error is that it enables us to have some backwards compatibility with OpenACC 2.0a. Unfortunately, the fortran FE doesn't make use of verify_oacc_routine_clauses because it parses/matches routines much earlier than c/c++, so that part of the patch is slightly more involved. I've also had to update a lot of test cases to make them conform to the new routine behavior. This patch has been committed to gomp-4_0-branch. Cesar
2017-05-01 Cesar Philippidis <ce...@codesourcery.com> gcc/fortran/ * gfortran.h (enum oacc_function): Add OACC_FUNCTION_AUTO. * openmp.c (gfc_oacc_routine_dims): Return OACC_FUNCTION_AUTO when no parallelism was detected. (gfc_match_oacc_routine): Emit a warning when the user doesn't supply a gang, worker, vector or seq clause to an OpenACC routine construct. gcc/ * omp-low.c (verify_oacc_routine_clauses): Emit a warning when the user doesn't supply a gang, worker, vector or seq clause to an OpenACC routine construct. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Update usage and expected output of OpenACC routines. * c-c++-common/goacc/Wparentheses-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Likewise. * c-c++-common/goacc/routine-1.c: Likewise. * 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: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * c-c++-common/goacc/routine-nohost-2.c: Likewise. * g++.dg/goacc/routine-1.C: Likewise. * g++.dg/goacc/routine-2.C: Likewise. * g++.dg/goacc/template.C: Likewise. * gfortran.dg/goacc/dtype-1.f95: Likewise. * gfortran.dg/goacc/pr71704.f90: Likewise. * gfortran.dg/goacc/pr72741-2.f: Likewise. * gfortran.dg/goacc/routine-6.f90: Likewise. * gfortran.dg/goacc/routine-8.f90: Likewise. * gfortran.dg/goacc/routine-9.f90: Likewise. * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise. * gfortran.dg/goacc/routine-without-clauses.f90: New test. libgomp/ * testsuite/libgomp.oacc-c++/pr71959-a.C: Adjust test case conform to OpenACC 2.5 routines. * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-default.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c: Likewise. * testsuite/libgomp.oacc-fortran/host_data-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/host_data-3.f: Likewise. * testsuite/libgomp.oacc-fortran/host_data-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-9.f90: Likewise. diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 7913ac7..3c762a8 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -316,7 +316,8 @@ enum oacc_function OACC_FUNCTION_GANG, OACC_FUNCTION_WORKER, OACC_FUNCTION_VECTOR, - OACC_FUNCTION_SEQ + OACC_FUNCTION_SEQ, + OACC_FUNCTION_AUTO }; /* Strings for all symbol attributes. We use these for dumping the diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 72c6669..88ccff2 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -2379,7 +2379,7 @@ static oacc_function gfc_oacc_routine_dims (gfc_omp_clauses *clauses) { int level = -1; - oacc_function ret = OACC_FUNCTION_SEQ; + oacc_function ret = OACC_FUNCTION_AUTO; if (clauses) { @@ -2401,7 +2401,10 @@ gfc_oacc_routine_dims (gfc_omp_clauses *clauses) ret = OACC_FUNCTION_VECTOR; } if (clauses->seq) - level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level); + { + level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level); + ret = OACC_FUNCTION_SEQ; + } if (mask != (mask & -mask)) ret = OACC_FUNCTION_NONE; @@ -2505,6 +2508,12 @@ gfc_match_oacc_routine (void) know of any potential duplicate routine directives. */ seen_error = true; } + else if (dims == OACC_FUNCTION_AUTO) + { + gfc_warning (0, "Expected one of %<gang%>, %<worker%>, %<vector%> or " + "%<seq%> clauses in !$ACC ROUTINE at %L", &old_loc); + dims = OACC_FUNCTION_SEQ; + } if (isym != NULL) { diff --git a/gcc/omp-low.c b/gcc/omp-low.c index cc209ba..88a1bb9 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -13239,8 +13239,10 @@ verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc, } if (c_level == NULL_TREE) { - /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a - implementation add an implicit "seq" clause. */ + /* OpenACC 2.5 expects the user to supply one parallelism clause. */ + warning_at (loc, 0, "expecting one of %<gang%>, %<worker%>, %<vector%> " + "or %<seq%> clauses"); + inform (loc, "assigning %<seq%> parallelism to this routine"); c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ); OMP_CLAUSE_CHAIN (c_level) = *clauses; *clauses = c_level; diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 1a33242..57eaa02 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -362,7 +362,7 @@ f_acc_data (void) } } -#pragma acc routine +#pragma acc routine seq void f_acc_loop (void) { @@ -436,7 +436,7 @@ f_acc_loop (void) } } -#pragma acc routine +#pragma acc routine seq void f_acc_routine (void) { diff --git a/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c b/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c index 08265b6..9d35ec5 100644 --- a/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c +++ b/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c @@ -4,9 +4,9 @@ int a, b, c; void bar (void); void baz (void); -#pragma acc routine +#pragma acc routine seq void bar2 (void); -#pragma acc routine +#pragma acc routine seq void baz2 (void); void diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 93a9111..1462095 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -57,7 +57,7 @@ f_acc_data (void) } } -#pragma acc routine +#pragma acc routine seq void f_acc_routine (void) { diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index 0b56661..a4ecfd3 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -21,15 +21,15 @@ void seq (void) } -#pragma acc routine +#pragma acc routine /* { dg-warning "expecting one of" } */ void bind_f_1 (void) { } -#pragma acc routine +#pragma acc routine /* { dg-warning "expecting one of" } */ extern void bind_f_1 (void); -#pragma acc routine (bind_f_1) +#pragma acc routine (bind_f_1)/* { dg-warning "expecting one of" } */ #pragma acc routine \ @@ -106,3 +106,16 @@ int main () return 0; } + +/* { dg-warning "expecting one of" "" { target *-*-* } 35 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 41 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 45 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 50 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 56 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 60 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 64 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 70 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 74 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 78 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 84 } */ +/* { dg-warning "expecting one of" "" { target *-*-* } 88 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c index debf6d7..65c7963 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -1,171 +1,171 @@ /* 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" } */ +#pragma acc routine seq 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" } */ +#pragma acc routine seq 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" }*/ +#pragma acc routine seq 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" } */ +#pragma acc routine seq 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" } */ +#pragma acc routine seq 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" } */ +#pragma acc routine seq 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" } */ +#pragma acc routine seq 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 (a_bind_f_1) seq #pragma acc routine \ - bind (a_bind_f_1) + bind (a_bind_f_1) seq 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" } */ +#pragma acc routine seq /* { 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" } */ +#pragma acc routine (a_bind_f_1_1) seq /* { 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") + bind ("a_bind_f_2") seq 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" } */ +#pragma acc routine seq /* { 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 (a_bind_f_2) seq /* { 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") + bind ("a_bind_f_2") seq 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" } */ +#pragma acc routine seq /* { 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" } */ +#pragma acc routine (a_bind_f_2_1) seq /* { 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 +#pragma acc routine seq extern void b_bind_f_1 (void); -#pragma acc routine +#pragma acc routine seq 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" } */ + bind (b_bind_f_1) seq /* { 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" } */ + bind (b_bind_f_1) seq /* { 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 +#pragma acc routine seq 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" } */ + bind ("b_bind_f_2") seq /* { 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" } */ + bind ("b_bind_f_2") seq /* { 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 +#pragma acc routine seq 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" } */ + bind ("b_bind_f_2") seq /* { 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" } */ + bind ("b_bind_f_2") seq /* { 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 +#pragma acc routine seq void c_bind_f_1a (void) { } -#pragma acc routine +#pragma acc routine seq extern void c_bind_f_1b (void); #pragma acc routine \ - bind (c_bind_f_1a) + bind (c_bind_f_1a) seq 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" } */ + bind (c_bind_f_1b) seq /* { 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" } */ + bind (c_bind_f_1b) seq /* { 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") + bind ("c_bind_f_2") seq 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" } */ + bind ("C_BIND_F_2") seq /* { 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" } */ + bind ("C_BIND_F_2") seq /* { 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") + bind ("c_bind_f_2") seq 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" } */ + bind ("C_BIND_F_2") seq /* { 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" } */ + bind ("C_BIND_F_2") seq /* { 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 f10651d..b7e8402 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -4,11 +4,11 @@ struct PC { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ }; void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } 11 } { dg-error ".#pragma. is not allowed here" "" { target c++ } 11 } */ ) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */ @@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c void PC2() { if (0) -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ ; } void PC3() { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ } /* "( name )" syntax. */ #pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */ -#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ -#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ -#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ -#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ -#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ +#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ +#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ +#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ +#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ +#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ #pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */ -#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */ +#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */ extern void R1(void); extern void R2(void); #pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */ @@ -49,84 +49,84 @@ extern void R2(void); /* "#pragma acc routine" not immediately followed by (a single) function declaration or definition. */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int a; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1 (void), fn1b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b, fn2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_, fn2_ (void), B_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3 (void), b2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c c; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d {} d; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1_2 (void), fn1b_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2, fn2_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2_, fn2_2_ (void), B_2_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3_2 (void), b2_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c_2 c_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d_2 {} d_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq int fn4 (void); int fn5a (void); int fn5b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn5a) -#pragma acc routine (fn5b) +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn5a) seq +#pragma acc routine (fn5b) seq int fn5 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */ -#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */ +#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */ int fn6 (void); #ifdef __cplusplus -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ namespace f {} namespace g {} -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ using namespace g; -#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ +#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ #endif /* __cplusplus */ -#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */ +#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */ -#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */ +#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */ /* Static assert. */ @@ -143,24 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11 #endif void f_static_assert(); /* Check that we already recognized "f_static_assert" as an OpenACC routine. */ -#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ +#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ /* __extension__ usage. */ -#pragma acc routine +#pragma acc routine seq __extension__ extern void ex1(); #pragma acc routine (ex1) worker /* { dg-error "has already been marked as an accelerator routine" } */ -#pragma acc routine +#pragma acc routine seq __extension__ __extension__ __extension__ __extension__ __extension__ 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" } */ +#pragma acc routine seq /* { 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 (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */ /* "#pragma acc routine" must be applied before. */ @@ -172,11 +172,11 @@ void Foo () Bar (); } -#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" } +#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" } #pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" } -#pragma acc routine (Baz) // { dg-error "not been declared" } +#pragma acc routine (Baz) seq // { dg-error "not been declared" } /* OpenACC declare. */ @@ -185,7 +185,7 @@ int vb1; /* { dg-error "directive for use" } */ extern int vb2; /* { dg-error "directive for use" } */ static int vb3; /* { dg-error "directive for use" } */ -#pragma acc routine +#pragma acc routine seq int func1 (int a) { @@ -196,7 +196,7 @@ func1 (int a) return vb3; } -#pragma acc routine +#pragma acc routine seq int func2 (int a) { @@ -214,7 +214,7 @@ extern int vb6; /* { dg-error "clause used in" } */ static int vb7; /* { dg-error "clause used in" } */ #pragma acc declare link (vb7) -#pragma acc routine +#pragma acc routine seq int func3 (int a) { @@ -231,7 +231,7 @@ extern int vb9; static int vb10; #pragma acc declare create (vb10) -#pragma acc routine +#pragma acc routine seq int func4 (int a) { @@ -249,7 +249,7 @@ extern int vb12; extern int vb13; #pragma acc declare device_resident (vb13) -#pragma acc routine +#pragma acc routine seq int func5 (int a) { @@ -260,7 +260,7 @@ func5 (int a) return vb13; } -#pragma acc routine +#pragma acc routine seq int func6 (int a) { 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 index 8f45499..e32d89d 100644 --- 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 @@ -428,20 +428,20 @@ extern void s_7 (void); 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 (g_8) seq /* { 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 (w_8) seq /* { 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" } */ +#pragma acc routine (v_8) seq /* { 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) /* { dg-warning "expecting one of" } */ #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) \ 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 index 1803997..7f26ef1 100644 --- 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 @@ -28,46 +28,46 @@ extern void v_1 (void); #pragma acc routine seq extern void s_1_1 (void); -#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_1_1) seq -#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_1_1) seq -#pragma acc routine +#pragma acc routine /* { dg-warning "expecting one of" } */ extern void s_1_2 (void); -#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_1_2) seq -#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) /* { dg-warning "expecting one of" } */ #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) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_2_1) seq -#pragma acc routine (s_2_1) +#pragma acc routine (s_2_1) /* { dg-warning "expecting one of" } */ #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) /* { dg-warning "expecting one of" } */ +#pragma acc routine (s_2_2) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_2_2) seq -#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) /* { dg-warning "expecting one of" } */ #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) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_3_1) seq -#pragma acc routine (s_3_1) +#pragma acc routine (s_3_1) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_3_1) seq -#pragma acc routine +#pragma acc routine seq void s_3_2 (void) { } -#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) /* { dg-warning "expecting one of" } */ #pragma acc routine (s_3_2) seq -#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) /* { dg-warning "expecting one of" } */ #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 index 17d6b03..800bcf6 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -3,46 +3,46 @@ /* { dg-additional-options "-fdump-tree-oaccdevlow" } */ -#pragma acc routine nohost +#pragma acc routine nohost seq int THREE(void) { return 3; } -#pragma acc routine (THREE) nohost +#pragma acc routine (THREE) nohost seq -#pragma acc routine nohost +#pragma acc routine nohost seq extern int THREE(void); -#pragma acc routine nohost +#pragma acc routine nohost seq extern void NOTHING(void); -#pragma acc routine (NOTHING) nohost +#pragma acc routine (NOTHING) nohost seq void NOTHING(void) { } -#pragma acc routine nohost +#pragma acc routine nohost seq extern void NOTHING(void); -#pragma acc routine (NOTHING) nohost +#pragma acc routine (NOTHING) nohost seq extern float ADD(float, float); -#pragma acc routine (ADD) nohost +#pragma acc routine (ADD) nohost seq float ADD(float x, float y) { return x + y; } -#pragma acc routine nohost +#pragma acc routine nohost seq extern float ADD(float, float); -#pragma acc routine (ADD) nohost +#pragma acc routine (ADD) nohost seq /* { 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 index 7402bfc..f172c38 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c @@ -1,41 +1,41 @@ /* Test invalid usage of the nohost clause for OpenACC routine directive. Exercising different variants for declaring routines. */ -#pragma acc routine +#pragma acc routine seq 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" } */ + nohost seq /* { 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" } */ + nohost seq /* { 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 +#pragma acc routine seq 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" } */ + nohost seq /* { 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" } */ + nohost seq /* { 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" } */ + nohost seq /* { 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) +#pragma acc routine (ADD_1) seq float ADD_1(float x, float y) { @@ -43,55 +43,55 @@ float ADD_1(float x, float 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" } */ + nohost seq /* { 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" } */ + nohost seq /* { 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 + nohost seq 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 (THREE_2) seq /* { 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" } */ +#pragma acc routine seq /* { 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 + nohost seq 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" } */ +#pragma acc routine (NOTHING_2) seq /* { 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" } */ +#pragma acc routine seq /* { 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" } */ +#pragma acc routine (NOTHING_2) seq /* { 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 + nohost seq 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" } */ +#pragma acc routine seq /* { 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" } */ +#pragma acc routine (ADD_2) seq /* { 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/g++.dg/goacc/routine-1.C b/gcc/testsuite/g++.dg/goacc/routine-1.C index a73a73d..9257324 100644 --- a/gcc/testsuite/g++.dg/goacc/routine-1.C +++ b/gcc/testsuite/g++.dg/goacc/routine-1.C @@ -4,10 +4,10 @@ namespace N { extern void foo1(); extern void foo2(); -#pragma acc routine (foo1) -#pragma acc routine +#pragma acc routine (foo1) seq +#pragma acc routine seq void foo3() { } } -#pragma acc routine (N::foo2) +#pragma acc routine (N::foo2) seq diff --git a/gcc/testsuite/g++.dg/goacc/routine-2.C b/gcc/testsuite/g++.dg/goacc/routine-2.C index 939dfba..846f388 100644 --- a/gcc/testsuite/g++.dg/goacc/routine-2.C +++ b/gcc/testsuite/g++.dg/goacc/routine-2.C @@ -2,7 +2,7 @@ template <typename T> extern T one_d(); -#pragma acc routine (one_d) nohost /* { dg-error "names a set of overloads" } */ +#pragma acc routine (one_d) nohost seq /* { dg-error "names a set of overloads" } */ template <typename T> T @@ -10,41 +10,41 @@ one() { return 1; } -#pragma acc routine (one) bind(one_d) /* { dg-error "names a set of overloads" } */ +#pragma acc routine (one) bind(one_d) seq /* { dg-error "names a set of overloads" } */ int incr (int); float incr (float); -#pragma acc routine (incr) /* { dg-error "names a set of overloads" } */ +#pragma acc routine (incr) seq /* { dg-error "names a set of overloads" } */ int sum (int, int); namespace foo { -#pragma acc routine (sum) +#pragma acc routine (sum) seq int sub (int, int); } -#pragma acc routine (foo::sub) +#pragma acc routine (foo::sub) seq /* It's strange to apply a routine directive to subset of overloaded functions, but that is permissible in OpenACC 2.x. */ int decr (int a); -#pragma acc routine +#pragma acc routine seq float decr (float a); /* Bind clause. */ -#pragma acc routine +#pragma acc routine seq float mult (float a, float b) { return a * b; } -#pragma acc routine bind(mult) /* { dg-error "bind identifier 'mult' is not compatible with function 'broken_mult1'" } */ +#pragma acc routine bind(mult) seq /* { dg-error "bind identifier 'mult' is not compatible with function 'broken_mult1'" } */ float broken_mult1 (int a, int b) { @@ -52,51 +52,51 @@ broken_mult1 (int a, int b) } /* This should result in a link error, but it's valid for a compile test. */ -#pragma acc routine bind("mult") +#pragma acc routine bind("mult") seq float broken_mult2 (float a, float b) { return a + b; } -#pragma acc routine (broken_mult2) bind("mult") -#pragma acc routine bind("mult") +#pragma acc routine (broken_mult2) bind("mult") seq +#pragma acc routine bind("mult") seq extern float broken_mult2 (float, float); -#pragma acc routine bind(sum2) /* { dg-error "'sum2' has not been declared" } */ +#pragma acc routine bind(sum2) seq /* { dg-error "'sum2' has not been declared" } */ int broken_mult3 (int a, int b); -#pragma acc routine bind(foo::sub) +#pragma acc routine bind(foo::sub) seq int broken_mult4 (int a, int b); -#pragma acc routine (broken_mult4) bind(foo::sub) -#pragma acc routine bind(foo::sub) +#pragma acc routine (broken_mult4) bind(foo::sub) seq +#pragma acc routine bind(foo::sub) seq extern int broken_mult4 (int a, int b); namespace bar { -#pragma acc routine bind (mult) +#pragma acc routine bind (mult) seq float working_mult (float a, float b) { return a * b; } -#pragma acc routine (working_mult) bind (mult) -#pragma acc routine bind (mult) +#pragma acc routine (working_mult) bind (mult) seq +#pragma acc routine bind (mult) seq float working_mult (float, float); } -#pragma acc routine +#pragma acc routine seq int div (int, int); -#pragma acc routine +#pragma acc routine seq float div (float, float); -#pragma acc routine bind (div) /* { dg-error "'div' names a set of overloads" } */ +#pragma acc routine bind (div) seq /* { dg-error "'div' names a set of overloads" } */ float my_div (float a, float b) { return a / b; } -#pragma acc routine bind (other_div) /* { dg-error "'other_div' has not been declared" } */ +#pragma acc routine bind (other_div) seq /* { dg-error "'other_div' has not been declared" } */ float my_div2 (float a, float b) { @@ -105,21 +105,21 @@ my_div2 (float a, float b) int div_var; -#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */ +#pragma acc routine bind (div_var) seq /* { dg-error "'div_var' does not refer to a function" } */ float my_div3 (float, float); -#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */ +#pragma acc routine bind (div_var) seq /* { dg-error "'div_var' does not refer to a function" } */ float my_div4 (float, float); -#pragma acc routine bind (%) /* { dg-error "expected identifier or character string literal" } */ +#pragma acc routine bind (%) seq /* { dg-error "expected identifier or character string literal" } */ float my_div5 (float, float); -#pragma acc routine bind ("") /* { dg-error "bind argument must not be an empty string" } */ +#pragma acc routine bind ("") seq /* { dg-error "bind argument must not be an empty string" } */ float my_div6 (float, float); struct astruct { - #pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ + #pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ int sum (int a, int b) { return a + b; @@ -128,7 +128,7 @@ struct astruct class aclass { - #pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ + #pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ int sum (int a, int b) { return a + b; diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index 056398d..4bc2596 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -1,4 +1,4 @@ -#pragma acc routine +#pragma acc routine seq template <typename T> T accDouble(int val) { diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 index 5919ae4..f24b60f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 @@ -108,63 +108,63 @@ end program dtype !! ACC ROUTINE: subroutine sr1 () - !$acc routine device_type (nvidia) gang + !$acc routine seq device_type (nvidia) gang end subroutine sr1 subroutine sr2 () - !$acc routine dtype (nvidia) worker + !$acc routine seq dtype (nvidia) worker end subroutine sr2 subroutine sr3 () - !$acc routine device_type (nvidia) vector + !$acc routine seq device_type (nvidia) vector end subroutine sr3 subroutine sr4 () - !$acc routine device_type (nvidia) seq + !$acc routine seq device_type (nvidia) seq end subroutine sr4 subroutine sr5 () - !$acc routine dtype (nvidia) bind (foo) + !$acc routine seq dtype (nvidia) bind (foo) end subroutine sr5 subroutine sr1a () - !$acc routine device_type (nvidia) gang device_type (*) seq + !$acc routine seq device_type (nvidia) gang device_type (*) seq end subroutine sr1a subroutine sr2a () - !$acc routine dtype (nvidia) worker dtype (*) seq + !$acc routine seq dtype (nvidia) worker dtype (*) seq end subroutine sr2a subroutine sr3a () - !$acc routine dtype (nvidia) vector device_type (*) seq + !$acc routine seq dtype (nvidia) vector device_type (*) seq end subroutine sr3a subroutine sr4a () - !$acc routine device_type (nvidia) seq device_type (*) worker + !$acc routine seq device_type (nvidia) seq device_type (*) worker end subroutine sr4a subroutine sr5a () - !$acc routine device_type (nvidia) bind (foo) dtype (*) seq + !$acc routine seq device_type (nvidia) bind (foo) dtype (*) seq end subroutine sr5a subroutine sr1b () - !$acc routine dtype (gpu) gang dtype (*) seq + !$acc routine seq dtype (gpu) gang dtype (*) seq end subroutine sr1b subroutine sr2b () - !$acc routine dtype (gpu) worker device_type (*) seq + !$acc routine seq dtype (gpu) worker device_type (*) seq end subroutine sr2b subroutine sr3b () - !$acc routine device_type (gpu) vector device_type (*) seq + !$acc routine seq device_type (gpu) vector device_type (*) seq end subroutine sr3b subroutine sr4b () - !$acc routine device_type (gpu) seq device_type (*) worker + !$acc routine seq device_type (gpu) seq device_type (*) worker end subroutine sr4b subroutine sr5b () - !$acc routine dtype (gpu) bind (foo) device_type (*) seq + !$acc routine seq dtype (gpu) bind (foo) device_type (*) seq end subroutine sr5b ! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) num_gangs\\(100\\) num_workers\\(100\\) vector_length\\(32\\) \\\]" 1 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 index 0235e85..92d0c71 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 @@ -2,7 +2,7 @@ ! { dg-do compile } real function f1 () -!$acc routine (f1) +!$acc routine (f1) seq f1 = 1 end diff --git a/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f b/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f index 5865144..0f0f7df 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f +++ b/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f @@ -1,5 +1,5 @@ SUBROUTINE v_1 -!$ACC ROUTINE +!$ACC ROUTINE SEQ !$ACC ROUTINE ! { dg-error "ACC ROUTINE already applied" } !$ACC ROUTINE GANG ! { dg-error "ACC ROUTINE already applied" } !$ACC ROUTINE SEQ ! { dg-error "ACC ROUTINE already applied" } @@ -13,7 +13,7 @@ !$ACC ROUTINE (g_1) GANG !$ACC ROUTINE (g_1) VECTOR ! { dg-error "ACC ROUTINE already applied" } !$ACC ROUTINE (g_1) SEQ ! { dg-error "ACC ROUTINE already applied" } -!$ACC ROUTINE (g_1) ! { dg-error "ACC ROUTINE already applied" } +!$ACC ROUTINE (g_1) WORKER! { dg-error "ACC ROUTINE already applied" } !$ACC ROUTINE (g_1) ! { dg-error "ACC ROUTINE already applied" } CALL v_1 diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 index ee43935..5c1f652 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 @@ -4,7 +4,7 @@ module m contains subroutine subr5 (x) implicit none - !$acc routine (subr5) + !$acc routine (subr5) seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -18,7 +18,7 @@ program main implicit none interface function subr6 (x) - !$acc routine (subr6) ! { dg-error "without list is allowed in interface" } + !$acc routine (subr6) seq ! { dg-error "without list is allowed in interface" } integer, intent (in) :: x integer :: subr6 end function subr6 @@ -26,13 +26,13 @@ program main integer, parameter :: n = 10 integer :: a(n), i external :: subr2 - !$acc routine (subr2) + !$acc routine (subr2) seq external :: R1, R2 - !$acc routine (R1 R2 R3) ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" } - !$acc routine (R1, R2, R3) ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" } - !$acc routine (R1) - !$acc routine (R2) + !$acc routine (R1 R2 R3) seq ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" } + !$acc routine (R1, R2, R3) seq ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" } + !$acc routine (R1) seq + !$acc routine (R2) seq !$acc parallel !$acc loop @@ -44,7 +44,7 @@ program main end program main subroutine subr1 (x) - !$acc routine + !$acc routine seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -63,7 +63,7 @@ subroutine subr2 (x) end subroutine subr2 subroutine subr3 (x) - !$acc routine (subr3) + !$acc routine (subr3) seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -73,7 +73,7 @@ subroutine subr3 (x) end subroutine subr3 subroutine subr4 (x) - !$acc routine (subr4) + !$acc routine (subr4) seq integer, intent(inout) :: x if (x < 1) then x = 1 diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 index c903915..beca43f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 @@ -4,7 +4,7 @@ program main interface function s_1 (a) integer a - !$acc routine + !$acc routine seq end function s_1 end interface @@ -18,7 +18,7 @@ program main interface function s_3 (a) integer a - !$acc routine (s_3) ! { dg-error "Only the ..ACC ROUTINE form without list is allowed in interface block" } + !$acc routine (s_3) seq ! { dg-error "Only the ..ACC ROUTINE form without list is allowed in interface block" } end function s_3 end interface diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 index 590e594..4027471 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 @@ -7,9 +7,9 @@ contains subroutine subr5 (x) implicit none integer extfunc - !$acc routine (subr5) - !$acc routine (extfunc) - !$acc routine (m1int) ! { dg-error "invalid function name" } + !$acc routine (subr5) seq + !$acc routine (extfunc) seq + !$acc routine (m1int) seq ! { dg-error "invalid function name" } integer, intent(inout) :: x if (x < 1) then x = 1 @@ -29,13 +29,13 @@ program main end interface integer, parameter :: n = 10 integer :: a(n), i - !$acc routine (subr1) ! { dg-error "invalid function name" } + !$acc routine (subr1) seq ! { dg-error "invalid function name" } external :: subr2 - !$acc routine (subr2) + !$acc routine (subr2) seq external :: R1, R2 - !$acc routine (R1) - !$acc routine (R2) + !$acc routine (R1) seq + !$acc routine (R2) seq !$acc parallel !$acc loop @@ -47,7 +47,7 @@ program main end program main subroutine subr1 (x) - !$acc routine + !$acc routine seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -66,7 +66,7 @@ subroutine subr2 (x) end subroutine subr2 subroutine subr3 (x) - !$acc routine (subr3) + !$acc routine (subr3) seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -76,7 +76,7 @@ subroutine subr3 (x) end subroutine subr3 subroutine subr4 (x) - !$acc routine (subr4) + !$acc routine (subr4) seq integer, intent(inout) :: x if (x < 1) then x = 1 @@ -86,7 +86,7 @@ subroutine subr4 (x) end subroutine subr4 subroutine subr10 (x) - !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" } + !$acc routine (subr10) seq 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-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 index 364a058..7c245ce 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 @@ -9,7 +9,7 @@ subroutine g_1 ! { dg-warning "region is gang partitioned but does not contain g end subroutine g_1 subroutine s_1_2a - !$acc routine + !$acc routine seq end subroutine s_1_2a subroutine s_1_2b @@ -17,7 +17,7 @@ subroutine s_1_2b end subroutine s_1_2b subroutine s_1_2c - !$acc routine (s_1_2c) + !$acc routine (s_1_2c) seq end subroutine s_1_2c subroutine s_1_2d @@ -27,7 +27,7 @@ end subroutine s_1_2d module s_2 contains subroutine s_2_1a - !$acc routine + !$acc routine seq end subroutine s_2_1a subroutine s_2_1b @@ -35,7 +35,7 @@ contains end subroutine s_2_1b subroutine s_2_1c - !$acc routine (s_2_1c) + !$acc routine (s_2_1c) seq end subroutine s_2_1c subroutine s_2_1d @@ -50,7 +50,7 @@ subroutine test interface function s_3_1a (a) integer a - !$acc routine + !$acc routine seq end function s_3_1a end interface diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90 new file mode 100644 index 0000000..570c8ed --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90 @@ -0,0 +1,36 @@ +! Test the OpenACC routine directive when it has no gang, worker +! vector, or seq partitioning clauses. + +! { dg-do compile } + +subroutine s1 + !$acc routine ! { dg-warning "Expected one of" } +end subroutine s1 + +integer function f1 () + !$acc routine ! { dg-warning "Expected one of" } +end function f1 + +module m +contains + subroutine s2 + !$acc routine ! { dg-warning "Expected one of" } + end subroutine s2 + + integer function f2 () + !$acc routine ! { dg-warning "Expected one of" } + end function f2 +end module m + +program t + implicit none + interface + subroutine s3 + !$acc routine ! { dg-warning "Expected one of" } + end subroutine s3 + + integer function f3 () + !$acc routine ! { dg-warning "Expected one of" } + end function f3 + end interface +end program t diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C index 9486512..82d6c22 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C +++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C @@ -8,13 +8,13 @@ struct Iter int *point () const asm("_ZNK4Iter5pointEv"); }; -#pragma acc routine +#pragma acc routine seq void Iter::ctor (int *cursor_) { cursor = cursor_; } -#pragma acc routine +#pragma acc routine seq int *Iter::point () const { return cursor; @@ -22,7 +22,7 @@ int *Iter::point () const void apply (int (*fn)(), Iter out) asm ("_ZN5Apply5applyEPFivE4Iter"); -#pragma acc routine +#pragma acc routine seq void apply (int (*fn)(), struct Iter out) { *out.point() = fn (); } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c index 2078a33..4f06ffd 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c @@ -7,7 +7,7 @@ float c[N]; #pragma acc declare device_resident (c) -#pragma acc routine +#pragma acc routine seq float subr2 (float a) { @@ -25,7 +25,7 @@ subr2 (float a) float b[N]; #pragma acc declare copyin (b) -#pragma acc routine +#pragma acc routine seq float subr1 (float a) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c index c3a2187..1b72bda 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c @@ -6,7 +6,7 @@ float *b; #pragma acc declare deviceptr (b) -#pragma acc routine +#pragma acc routine seq float * subr2 (void) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c index 36bf0eb..385b710 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c @@ -6,7 +6,7 @@ float b; #pragma acc declare create (b) -#pragma acc routine +#pragma acc routine seq int func (int a) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c index fe843ec..9cb4665 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c @@ -8,7 +8,7 @@ #include <cuda_runtime_api.h> #include <cublas_v2.h> -#pragma acc routine +#pragma acc routine seq void saxpy (int n, float a, float *x, float *y) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h index 55de04b..36e8497 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h @@ -3,7 +3,7 @@ #include <string.h> #include <stdio.h> -#pragma acc routine +#pragma acc routine seq static int __attribute__ ((noinline)) coord () { int res = 0; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c index f62daf0..d3a5858 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c @@ -1122,7 +1122,7 @@ void t27() /* Check if worker-single variables get broadcastd to vectors. */ -#pragma acc routine +#pragma acc routine seq float t28_routine () { return 2.71; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h index 5691b7e..1934a2b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h @@ -2,7 +2,7 @@ #define VARS int a[1500]; float b[10][15][10]; -#pragma acc routine +#pragma acc routine seq __attribute__((noreturn)) void noreturn (void) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c index 0f70e26..e45e930 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c @@ -8,7 +8,7 @@ #include <stdlib.h> -#pragma acc routine +#pragma acc routine seq TEMPLATE RETURN_1 fact(TYPE n) RETURN_2 { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c index 2cdd6bf..07b9551 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c @@ -6,7 +6,7 @@ #include <stdlib.h> -#pragma acc routine nohost +#pragma acc routine nohost seq int foo (int n) { 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 index b991bb1..ec4a4a3 100644 --- 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 @@ -13,7 +13,7 @@ /* "MINUS_TWO" is the device variant for function "TWO". Similar for "THREE", and "FOUR". Exercising different variants for declaring routines. */ -#pragma acc routine nohost +#pragma acc routine nohost seq extern int MINUS_TWO(void); int MINUS_TWO(void) @@ -24,7 +24,7 @@ int MINUS_TWO(void) } extern int TWO(void); -#pragma acc routine (TWO) bind(MINUS_TWO) +#pragma acc routine (TWO) bind(MINUS_TWO) seq int TWO(void) { @@ -34,7 +34,7 @@ int TWO(void) } -#pragma acc routine nohost +#pragma acc routine nohost seq int MINUS_THREE(void) { if (!acc_on_device(acc_device_not_host)) @@ -42,7 +42,7 @@ int MINUS_THREE(void) return -3; } -#pragma acc routine bind(MINUS_THREE) +#pragma acc routine bind(MINUS_THREE) seq extern int THREE(void); int THREE(void) @@ -55,7 +55,7 @@ int THREE(void) /* Due to using a string in the bind clause, we don't need "MINUS_FOUR" in scope here. */ -#pragma acc routine bind("MINUS_FOUR") +#pragma acc routine bind("MINUS_FOUR") seq int FOUR(void) { if (acc_on_device(acc_device_not_host)) @@ -64,7 +64,7 @@ int FOUR(void) } extern int MINUS_FOUR(void); -#pragma acc routine (MINUS_FOUR) nohost +#pragma acc routine (MINUS_FOUR) nohost seq int MINUS_FOUR(void) { diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90 index ff09218..4cdaa0c 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90 @@ -81,7 +81,7 @@ subroutine saxpy (nn, aa, xx, yy) integer :: nn real*4 :: aa, xx(nn), yy(nn) integer i - !$acc routine + !$acc routine seq do i = 1, nn yy(i) = yy(i) + aa * xx(i) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f b/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f index 05ed949..fe0b904 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f @@ -67,7 +67,7 @@ integer :: nn real*4 :: aa, xx(nn), yy(nn) integer i -!$acc routine +!$acc routine seq do i = 1, nn yy(i) = yy(i) + aa * xx(i) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90 index 6e379b5..e192c59 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90 @@ -21,7 +21,7 @@ contains integer :: nn real*4 :: aa, xx(nn), yy(nn) integer i - !$acc routine + !$acc routine seq do i = 1, nn yy(i) = yy(i) + aa * xx(i) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90 index 3390515..3b85391 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90 @@ -3,7 +3,7 @@ interface recursive function fact (x) - !$acc routine + !$acc routine seq integer, intent(in) :: x integer :: fact end function fact @@ -21,7 +21,7 @@ end do end recursive function fact (x) result (res) - !$acc routine + !$acc routine seq integer, intent(in) :: x integer :: res if (x < 1) then diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90 index 3d418b6..902ab69 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90 @@ -4,7 +4,7 @@ module m1 contains recursive function fact (x) result (res) - !$acc routine + !$acc routine seq integer, intent(in) :: x integer :: res if (x < 1) then diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90 index d233a63..c3f6edd 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90 @@ -4,7 +4,7 @@ integer, parameter :: n = 10 integer :: a(n), i integer, external :: fact - !$acc routine (fact) + !$acc routine (fact) seq !$acc parallel !$acc loop do i = 1, n @@ -16,7 +16,7 @@ end do end recursive function fact (x) result (res) - !$acc routine + !$acc routine seq integer, intent(in) :: x integer :: res if (x < 1) then diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90 index 3e5fb09..7e763b8 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90 @@ -17,7 +17,7 @@ end do end subroutine incr (x) - !$acc routine + !$acc routine seq integer, intent(inout) :: x x = x + 1 end subroutine incr diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 index 956da8e..8c5d03a 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 @@ -15,7 +15,7 @@ program main contains function func (n) result (rc) - !$acc routine + !$acc routine seq integer, intent (in) :: n integer :: rc diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90 index 95d1a13..6f05e28 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90 @@ -6,7 +6,7 @@ program main integer, parameter :: n = 10 integer :: a(n), i integer, external :: fact - !$acc routine (fact) + !$acc routine (fact) seq !$acc parallel !$acc loop do i = 1, n @@ -20,7 +20,7 @@ end program main recursive function fact (x) result (res) implicit none - !$acc routine (fact) + !$acc routine (fact) seq integer, intent(in) :: x integer :: res if (x < 1) then