diff mbox series

Handle 'omp declare target' attribute set for both OpenACC and OpenMP 'target' [PR89433, PR93465] (was: [committed] PR89433 "Repeated use of the OpenACC 'routine' directive")

Message ID 87k13zly41.fsf@euler.schwinge.homeip.net
State New
Headers show
Series Handle 'omp declare target' attribute set for both OpenACC and OpenMP 'target' [PR89433, PR93465] (was: [committed] PR89433 "Repeated use of the OpenACC 'routine' directive") | expand

Commit Message

Thomas Schwinge March 4, 2020, 7:27 p.m. UTC
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi!

On 2019-05-17T21:16:57+0200, I wrote:
> Now committed to trunk in [...]
> r271345 "[PR89433] Repeated use of the C/C++ OpenACC 'routine'
> directive"

> --- a/gcc/omp-general.c
> +++ b/gcc/omp-general.c

> @@ -610,11 +610,14 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
>  
>  /* Verify OpenACC routine clauses.
>  
> +   Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
> +   if it has already been marked in compatible way, and -1 if incompatible.
>     Upon returning, the chain of clauses will contain exactly one clause
>     specifying the level of parallelism.  */

> +  tree attr
> +    = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
> +  if (attr != NULL_TREE)
> +    {
> +      /* If a "#pragma acc routine" has already been applied, just verify
> +	 this one for compatibility.  */
> +      /* Collect previous directive's clauses.  */
> +      tree c_level_p = NULL_TREE;
> +      for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
> +	switch (OMP_CLAUSE_CODE (c))
> +	  {
> +	  case OMP_CLAUSE_GANG:
> +	  case OMP_CLAUSE_WORKER:
> +	  case OMP_CLAUSE_VECTOR:
> +	  case OMP_CLAUSE_SEQ:
> +	    gcc_checking_assert (c_level_p == NULL_TREE);
> +	    c_level_p = c;
> +	    break;
> +	  default:
> +	    gcc_unreachable ();
> +	  }
> +      gcc_checking_assert (c_level_p != NULL_TREE);

As documented in <https://gcc.gnu.org/PR93465>, this triggers an ICE if
the 'omp declare target' attribute had already been set for '#pragma omp
declare target'.  OK to deal with this situation as in the patch
attached, "Handle 'omp declare target' attribute set for both OpenACC and
OpenMP 'target' [PR89433, PR93465]"?  If approving this patch, please
respond with "Reviewed-by: NAME <EMAIL>" so that your effort will be
recorded in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.

That's probably not worth backporting (a variant of this) to release
branches -- where also other such mixed OpenACC/OpenMP code is silently
accepted, with unclear semantics.


Grüße
 Thomas

Comments

Jakub Jelinek March 4, 2020, 8:07 p.m. UTC | #1
On Wed, Mar 04, 2020 at 08:27:10PM +0100, Thomas Schwinge wrote:
> ... which as of PR89433 commit b48f44bf77a39fefc238a16cf1225c6464c82406 causes
> an ICE.  Not sure if this is actually supposed to be valid or invalid code.
> Until the interactions between OpenACC and OpenMP 'target' get defined
> properly, make this a compile-time error.
> 
> 	gcc/
> 	PR middle-end/89433
> 	PR middle-end/93465
> 	* omp-general.c (oacc_verify_routine_clauses): Diagnose if
> 	"#pragma omp declare target" has also been applied.
> 	gcc/testsuite/
> 	PR middle-end/89433
> 	PR middle-end/93465
> 	* c-c++-common/goacc-gomp/pr93465-1.c: New file.

Ok for trunk.

	Jakub
Thomas Schwinge April 10, 2020, 2:14 p.m. UTC | #2
Hi!

On 2020-03-04T20:07:46+0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Mar 04, 2020 at 08:27:10PM +0100, Thomas Schwinge wrote:
>> ... which as of PR89433 commit b48f44bf77a39fefc238a16cf1225c6464c82406 causes
>> an ICE.  Not sure if this is actually supposed to be valid or invalid code.
>> Until the interactions between OpenACC and OpenMP 'target' get defined
>> properly, make this a compile-time error.
>>
>>      gcc/
>>      PR middle-end/89433
>>      PR middle-end/93465
>>      * omp-general.c (oacc_verify_routine_clauses): Diagnose if
>>      "#pragma omp declare target" has also been applied.
>>      gcc/testsuite/
>>      PR middle-end/89433
>>      PR middle-end/93465
>>      * c-c++-common/goacc-gomp/pr93465-1.c: New file.
>
> Ok for trunk.

Thanks for the review, and sorry for the delay; now pushed to master
branch in commit ff3f862b451496dd4afbe2dbfae82afab59a42c6 "Handle 'omp
declare target' attribute set for both OpenACC and OpenMP 'target'
[PR89433, PR93465]", see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
diff mbox series

Patch

From 3f8e048f5f8d1eabde642c1c146114027bb44e79 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 4 Mar 2020 17:58:33 +0100
Subject: [PATCH] Handle 'omp declare target' attribute set for both OpenACC
 and OpenMP 'target' [PR89433, PR93465]

... which as of PR89433 commit b48f44bf77a39fefc238a16cf1225c6464c82406 causes
an ICE.  Not sure if this is actually supposed to be valid or invalid code.
Until the interactions between OpenACC and OpenMP 'target' get defined
properly, make this a compile-time error.

	gcc/
	PR middle-end/89433
	PR middle-end/93465
	* omp-general.c (oacc_verify_routine_clauses): Diagnose if
	"#pragma omp declare target" has also been applied.
	gcc/testsuite/
	PR middle-end/89433
	PR middle-end/93465
	* c-c++-common/goacc-gomp/pr93465-1.c: New file.
---
 gcc/omp-general.c                             | 13 +++++
 .../c-c++-common/goacc-gomp/pr93465-1.c       | 56 +++++++++++++++++++
 2 files changed, 69 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/goacc-gomp/pr93465-1.c

diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index f107f4c050f1..49023f42c473 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -1776,6 +1776,19 @@  oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
     = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
   if (attr != NULL_TREE)
     {
+      /* Diagnose if "#pragma omp declare target" has also been applied.  */
+      if (TREE_VALUE (attr) == NULL_TREE)
+	{
+	  /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
+	     OpenACC and OpenMP 'target' are not clear.  */
+	  error_at (loc,
+		    "cannot apply %<%s%> to %qD, which has also been"
+		    " marked with an OpenMP 'declare target' directive",
+		    routine_str, fndecl);
+	  /* Incompatible.  */
+	  return -1;
+	}
+
       /* If a "#pragma acc routine" has already been applied, just verify
 	 this one for compatibility.  */
       /* Collect previous directive's clauses.  */
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/pr93465-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/pr93465-1.c
new file mode 100644
index 000000000000..c8b9135d9973
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/pr93465-1.c
@@ -0,0 +1,56 @@ 
+#pragma omp declare target
+#pragma acc routine seq /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f1\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+void f1 (void) {}
+#pragma omp end declare target
+
+#pragma omp declare target
+void f1 (void);
+
+#pragma acc routine seq /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f1\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+void f1 (void);
+
+
+
+#pragma omp declare target
+#pragma acc routine /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f2\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+extern void f2 (void);
+#pragma omp end declare target
+
+#pragma omp declare target
+extern void f2 (void);
+#pragma omp end declare target
+
+#pragma acc routine gang /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f2\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+extern void f2 (void);
+
+
+#pragma omp declare target
+#pragma acc routine gang /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f3\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+void f3 (void);
+#pragma omp end declare target
+
+#pragma omp declare target
+void f3 (void) {}
+#pragma omp end declare target
+
+#pragma acc routine (f3) gang /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f3\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+
+
+/* Surprisingly, this diagnosis also works for '#pragma acc routine' first,
+   followed by '#pragma omp declare target'; the latter gets applied first.  */
+
+
+#pragma acc routine /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f4\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+extern void f4 (void);
+
+#pragma omp declare target
+extern void f4 (void);
+#pragma omp end declare target
+
+
+#pragma acc routine gang /* { dg-error "cannot apply '#pragma acc routine' to '\(void \)?f5\(\\(\\)\)?', which has also been marked with an OpenMP 'declare target' directive" } */
+void f5 (void) {}
+
+#pragma omp declare target
+extern void f5 (void);
+#pragma omp end declare target
-- 
2.17.1