Hi!
On Mon, 01 Aug 2016 17:21:37 +0200, I wrote:
> Some checking of OpenACC clauses currently done in the front ends should
> be moved later, and be unified. (Also, I suppose, for supporting of the
> device_type clause, such checking actually *must* be moved later, into
> the oaccdevlow pass, or similar.) Here is a first preparatory patch. OK
> for trunk?
>
> commit e02a9b65c505b404f8d985b0ec6ccb99d73515d3
> Author: Thomas Schwinge <tho...@codesourcery.com>
> Date: Wed Jul 27 15:54:38 2016 +0200
>
> Use verify_oacc_routine_clauses for C/C++
Here is a Fortran patch. This depends on other Fortran patches in flight
(such as Cesar's), and on PR72741 "Fortran OpenACC routine directive
doesn't properly handle clauses specifying the level of parallelism" be
resolved, and thereabouts, but I'm posting it anyway, in case anyone has
any review comments already. I suppose, to begin with, the call of
gfc_oacc_routine_dims will move later into the Fortran front end
pipeline, to the point then function declarations' attributes are set, or
similar. Also, as discussed already, the Fortran front end currently is
very "forgetful" in regards to OpenACC/OpenMP clauses' specific location
information, so we're not able at present to produce diagnostics with
precise location information.
commit 6480b966af617e61b35b59bb089dd009064743e5
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri Jul 29 13:12:49 2016 +0200
Use verify_oacc_routine_clauses for Fortran
gcc/fortran/
* openmp.c: Include "trans-stmt.h".
(gfc_oacc_routine_dims): Move function...
* trans-openmp.c: ... here, and re-implement it. Adjust all
users.
* trans-stmt.h (gfc_oacc_routine_dims): New prototype.
gcc/testsuite/
* gfortran.dg/goacc/pr72741.f90: Update.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f: New file.
---
gcc/fortran/openmp.c | 47 +---------
gcc/fortran/trans-openmp.c | 67 +++++++++++++
gcc/fortran/trans-stmt.h | 7 +-
gcc/testsuite/gfortran.dg/goacc/pr72741.f90 | 10 +-
.../goacc/routine-level-of-parallelism-1.f | 104 +++++++++++++++++++++
5 files changed, 181 insertions(+), 54 deletions(-)
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index f7a907d..b0a10a8 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see
#include "parse.h"
#include "diagnostic.h"
#include "gomp-constants.h"
+#include "trans-stmt.h"
/* Match an end of OpenMP directive. End of OpenMP directive is optional
whitespace, followed by '\n' or comment '!'. */
@@ -1714,44 +1715,6 @@ gfc_match_oacc_cache (void)
return MATCH_YES;
}
-/* Determine the loop level for a routine. Returns OACC_FUNCTION_NONE if
- any error is detected. */
-
-static oacc_function
-gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
-{
- int level = -1;
- oacc_function ret = OACC_FUNCTION_SEQ;
-
- if (clauses)
- {
- unsigned mask = 0;
-
- if (clauses->gang)
- {
- level = GOMP_DIM_GANG, mask |= GOMP_DIM_MASK (level);
- ret = OACC_FUNCTION_GANG;
- }
- if (clauses->worker)
- {
- level = GOMP_DIM_WORKER, mask |= GOMP_DIM_MASK (level);
- ret = OACC_FUNCTION_WORKER;
- }
- if (clauses->vector)
- {
- level = GOMP_DIM_VECTOR, mask |= GOMP_DIM_MASK (level);
- ret = OACC_FUNCTION_VECTOR;
- }
- if (clauses->seq)
- level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
-
- if (mask != (mask & -mask))
- ret = OACC_FUNCTION_NONE;
- }
-
- return ret;
-}
-
match
gfc_match_oacc_routine (void)
{
@@ -1828,13 +1791,7 @@ gfc_match_oacc_routine (void)
!= MATCH_YES))
return MATCH_ERROR;
- dims = gfc_oacc_routine_dims (c);
- if (dims == OACC_FUNCTION_NONE)
- {
- gfc_error ("Multiple loop axes specified for routine %C");
- gfc_current_locus = old_loc;
- return MATCH_ERROR;
- }
+ dims = gfc_oacc_routine_dims (c, old_loc);
if (isym != NULL)
/* There is nothing to do for intrinsic procedures. */
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 0d646ed..254732c 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -4570,3 +4570,70 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns)
DECL_ATTRIBUTES (fndecl) = c;
}
}
+
+/* Determine and verify the level of parallelism for an OpenACC routine. */
+
+oacc_function
+gfc_oacc_routine_dims (gfc_omp_clauses *clauses, locus location)
+{
+ /* This is implemented in terms of OMP_CLAUSE trees, so that we can use the
+ generic functions for checking validity. This has a little bit of
+ overhead, but as the number of clauses on OpenACC routine directives as
+ well as the number of OpenACC routine directives will both be rather
+ small, this is acceptable. */
+ tree clauses_t = NULL_TREE;
+ /* We don't have specific location information available for the individual
+ clauses... */
+ location_t loc = location.lb->location;
+ if (clauses)
+ {
+ if (clauses->gang)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_GANG);
+ OMP_CLAUSE_CHAIN (c) = clauses_t;
+ clauses_t = c;
+ }
+ if (clauses->worker)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_WORKER);
+ OMP_CLAUSE_CHAIN (c) = clauses_t;
+ clauses_t = c;
+ }
+ if (clauses->vector)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_VECTOR);
+ OMP_CLAUSE_CHAIN (c) = clauses_t;
+ clauses_t = c;
+ }
+ if (clauses->seq)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+ OMP_CLAUSE_CHAIN (c) = clauses_t;
+ clauses_t = c;
+ }
+ }
+ verify_oacc_routine_clauses (&clauses_t, loc);
+
+ gcc_checking_assert (clauses_t != NULL_TREE
+ && OMP_CLAUSE_CHAIN (clauses_t) == NULL_TREE);
+ oacc_function ret;
+ switch (OMP_CLAUSE_CODE (clauses_t))
+ {
+ case OMP_CLAUSE_GANG:
+ ret = OACC_FUNCTION_GANG;
+ break;
+ case OMP_CLAUSE_WORKER:
+ ret = OACC_FUNCTION_WORKER;
+ break;
+ case OMP_CLAUSE_VECTOR:
+ ret = OACC_FUNCTION_VECTOR;
+ break;
+ case OMP_CLAUSE_SEQ:
+ ret = OACC_FUNCTION_SEQ;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ return ret;
+}
diff --git gcc/fortran/trans-stmt.h gcc/fortran/trans-stmt.h
index f9c8e74..1e50298 100644
--- gcc/fortran/trans-stmt.h
+++ gcc/fortran/trans-stmt.h
@@ -63,12 +63,11 @@ tree gfc_trans_deallocate (gfc_code *);
tree gfc_trans_deallocate_array (tree);
/* trans-openmp.c */
+tree gfc_trans_oacc_declare (gfc_namespace *);
+tree gfc_trans_oacc_directive (gfc_code *);
tree gfc_trans_omp_directive (gfc_code *);
void gfc_trans_omp_declare_simd (gfc_namespace *);
-
-/* trans-openacc.c */
-tree gfc_trans_oacc_directive (gfc_code *);
-tree gfc_trans_oacc_declare (gfc_namespace *);
+oacc_function gfc_oacc_routine_dims (gfc_omp_clauses *, locus);
/* trans-io.c */
tree gfc_trans_open (gfc_code *);
diff --git gcc/testsuite/gfortran.dg/goacc/pr72741.f90
gcc/testsuite/gfortran.dg/goacc/pr72741.f90
index cf89727..a5cb218 100644
--- gcc/testsuite/gfortran.dg/goacc/pr72741.f90
+++ gcc/testsuite/gfortran.dg/goacc/pr72741.f90
@@ -1,12 +1,12 @@
SUBROUTINE v_1
- !$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE VECTOR WORKER ! { dg-error ".worker. specifies a conflicting
level of parallelism" }
END SUBROUTINE v_1
SUBROUTINE sub_1
IMPLICIT NONE
EXTERNAL :: g_1
- !$ACC ROUTINE (g_1) GANG WORKER ! { dg-error "Multiple loop axes" }
- !$ACC ROUTINE (ABORT) SEQ VECTOR ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (g_1) GANG WORKER ! { dg-error ".gang. specifies a conflicting
level of parallelism" }
+ !$ACC ROUTINE (ABORT) SEQ VECTOR ! { dg-error ".vector. specifies a
conflicting level of parallelism" }
CALL v_1
CALL g_1
@@ -16,8 +16,8 @@ END SUBROUTINE sub_1
MODULE m_w_1
IMPLICIT NONE
EXTERNAL :: w_1
- !$ACC ROUTINE (w_1) WORKER SEQ ! { dg-error "Multiple loop axes" }
- !$ACC ROUTINE (ABORT) VECTOR GANG ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (w_1) WORKER SEQ ! { dg-error ".worker. specifies a
conflicting level of parallelism" }
+ !$ACC ROUTINE (ABORT) VECTOR GANG ! { dg-error ".gang. specifies a
conflicting level of parallelism" }
CONTAINS
SUBROUTINE sub_2
diff --git gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f
gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f
new file mode 100644
index 0000000..e2f3c61
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f
@@ -0,0 +1,104 @@
+! Test various aspects of clauses specifying incompatible levels of
+! parallelism with the OpenACC routine directive. The C/C++ counterpart
+! is ../../c-c++-common/goacc/routine-level-of-parallelism-1.c.
+
+! We currently don't have precise location information for these, so can't
+! match them properly.
+! { dg-prune-output "specifies a conflicting level of parallelism" }
+
+ MODULE m_g_1
+ EXTERNAL :: g_1
+!$ACC ROUTINE (g_1) GANG GANG ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+ SUBROUTINE w_1
+!$ACC ROUTINE WORKER WORKER ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+ SUBROUTINE v_1
+!$ACC ROUTINE VECTOR VECTOR ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+ MODULE m_s_1
+ EXTERNAL :: s_1
+!$ACC ROUTINE (s_1) SEQ SEQ ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+
+ SUBROUTINE g_2
+!$ACC ROUTINE GANG GANG GANG ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+ MODULE m_w_2
+ EXTERNAL :: w_2
+!$ACC ROUTINE (w_2) WORKER WORKER WORKER ! { dg-error "Unclassifiable OpenACC
directive" }
+ END
+
+ MODULE m_v_2
+ EXTERNAL :: v_2
+!$ACC ROUTINE (v_2) VECTOR VECTOR VECTOR ! { dg-error "Unclassifiable OpenACC
directive" }
+ END
+
+ SUBROUTINE s_2
+!$ACC ROUTINE SEQ SEQ SEQ ! { dg-error "Unclassifiable OpenACC directive" }
+ END
+
+
+ SUBROUTINE g_3
+!$ACC ROUTINE
+!$ACC& GANG
+!$ACC& WORKER ! { dg-error ".worker. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+ END
+
+ MODULE m_g_3
+ EXTERNAL :: g_3
+!$ACC ROUTINE (g_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*g_3." "" { xfail *-*-* } }
+!$ACC& GANG
+!$ACC& SEQ ! { dg-error ".seq. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+
+!$ACC ROUTINE (g_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*g_3." "" { xfail *-*-* } }
+!$ACC& GANG
+!$ACC& VECTOR ! { dg-error ".vector. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+ END
+
+ MODULE m_w_3
+ EXTERNAL :: w_3
+!$ACC ROUTINE (w_3)
+!$ACC& WORKER
+!$ACC& VECTOR ! { dg-error ".vector. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+!$ACC ROUTINE (w_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*w_3." "" { xfail *-*-* } }
+!$ACC& WORKER
+!$ACC& GANG ! { dg-error ".gang. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+!$ACC ROUTINE (w_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*w_3." "" { xfail *-*-* } }
+!$ACC& WORKER
+!$ACC& SEQ ! { dg-error ".seq. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+ END
+
+ SUBROUTINE v_3
+!$ACC ROUTINE (v_3)
+!$ACC& VECTOR
+!$ACC& SEQ ! { dg-error ".seq. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+ END
+
+ MODULE m_v_3
+ EXTERNAL :: v_3
+!$ACC ROUTINE (v_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*v_3." "" { xfail *-*-* } }
+!$ACC& VECTOR
+!$ACC& WORKER ! { dg-error ".worker. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+!$ACC ROUTINE (v_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*v_3." "" { xfail *-*-* } }
+!$ACC& VECTOR
+!$ACC& GANG ! { dg-error ".gang. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+ END
+
+ MODULE m_s_3
+ EXTERNAL :: s_3
+!$ACC ROUTINE (s_3)
+!$ACC& SEQ
+!$ACC& GANG ! { dg-error ".gang. specifies a conflicting level of parallelism"
"" { xfail *-*-* } }
+!$ACC ROUTINE (s_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*s_3." "" { xfail *-*-* } }
+!$ACC& SEQ
+!$ACC& VECTOR ! { dg-error ".vector. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+!$ACC ROUTINE (s_3) ! { dg-error ".#pragma acc routine. already applied to
.\[void \]*s_3." "" { xfail *-*-* } }
+!$ACC& SEQ
+!$ACC& WORKER ! { dg-error ".worker. specifies a conflicting level of
parallelism" "" { xfail *-*-* } }
+ END
Grüße
Thomas
