diff mbox

Fix PR78027

Message ID bdd55ec0-ba12-4413-e79d-1b051533ea1b@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Dec. 8, 2016, 11:05 p.m. UTC
PR78027 was classified as a fortran bug, but the underlying problem
turned out to be more generic. Basically, the IPA-ICF pass usually
ignores functions with omp decl attributes as candidates for aliasing.
Usually that works for OpenACC, because offloaded functions generated by
OpenACC PARALLEL or KERNELS regions usually have an 'omp target
entrypoint' attribute. However that is not the case in two situations:

1. ACC ROUTINES do not have an 'omp target entrypoint' attribute.

2. Nested offloaded regions inside 'omp declare target' functions do not
have 'omp target entrypoint' attributes either.

The solution I chose for this problem is to teach the IPA-ICF pass to
ignore function with 'oacc *' decl attributes. Arguably 2) should be a
parser error when an OpenACC offloaded region is embedded within an
OpenMP offloaded function. The LTO linker does report an error for nvptx
targets, but not the host. With that in mind, this patch is still
necessary for case 1.

Is this OK for trunk?

Cesar

Comments

Steve Kargl Dec. 8, 2016, 11:20 p.m. UTC | #1
On Thu, Dec 08, 2016 at 03:05:10PM -0800, Cesar Philippidis wrote:

> 2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

> 

> 	PR fortran/78027

> 

> 	gcc/

> 	* ipa-icf.c (sem_function::parse): Don't process functions with

> 	oacc decl attributes, as they may be OpenACC routines.

> 

> 	gcc/testsuite/

> 	* c-c++-common/goacc/acc-icf.c: New test.

> 	* gfortran.dg/goacc/pr78027.f90: New test.

> 


The Fortran testcase is OK.  I probably don't have the
necessary background to review the ipa-icf.c code changes.

-- 
Steve
Jeff Law Dec. 9, 2016, 12:37 a.m. UTC | #2
On 12/08/2016 04:05 PM, Cesar Philippidis wrote:
> PR78027 was classified as a fortran bug, but the underlying problem

> turned out to be more generic. Basically, the IPA-ICF pass usually

> ignores functions with omp decl attributes as candidates for aliasing.

> Usually that works for OpenACC, because offloaded functions generated by

> OpenACC PARALLEL or KERNELS regions usually have an 'omp target

> entrypoint' attribute. However that is not the case in two situations:

>

> 1. ACC ROUTINES do not have an 'omp target entrypoint' attribute.

>

> 2. Nested offloaded regions inside 'omp declare target' functions do not

> have 'omp target entrypoint' attributes either.

>

> The solution I chose for this problem is to teach the IPA-ICF pass to

> ignore function with 'oacc *' decl attributes. Arguably 2) should be a

> parser error when an OpenACC offloaded region is embedded within an

> OpenMP offloaded function. The LTO linker does report an error for nvptx

> targets, but not the host. With that in mind, this patch is still

> necessary for case 1.

>

> Is this OK for trunk?

>

> Cesar

>

>

> trunk-pr78027.diff

>

>

> 2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

>

> 	PR fortran/78027

>

> 	gcc/

> 	* ipa-icf.c (sem_function::parse): Don't process functions with

> 	oacc decl attributes, as they may be OpenACC routines.

>

> 	gcc/testsuite/

> 	* c-c++-common/goacc/acc-icf.c: New test.

> 	* gfortran.dg/goacc/pr78027.f90: New test.

This follows the same approach as we do for openmp.  This is fine for 
the trunk.

Jeff
Cesar Philippidis Dec. 9, 2016, 2:05 p.m. UTC | #3
On 12/08/2016 04:37 PM, Jeff Law wrote:
> On 12/08/2016 04:05 PM, Cesar Philippidis wrote:

>> PR78027 was classified as a fortran bug, but the underlying problem

>> turned out to be more generic. Basically, the IPA-ICF pass usually

>> ignores functions with omp decl attributes as candidates for aliasing.

>> Usually that works for OpenACC, because offloaded functions generated by

>> OpenACC PARALLEL or KERNELS regions usually have an 'omp target

>> entrypoint' attribute. However that is not the case in two situations:

>>

>> 1. ACC ROUTINES do not have an 'omp target entrypoint' attribute.

>>

>> 2. Nested offloaded regions inside 'omp declare target' functions do not

>> have 'omp target entrypoint' attributes either.

>>

>> The solution I chose for this problem is to teach the IPA-ICF pass to

>> ignore function with 'oacc *' decl attributes. Arguably 2) should be a

>> parser error when an OpenACC offloaded region is embedded within an

>> OpenMP offloaded function. The LTO linker does report an error for nvptx

>> targets, but not the host. With that in mind, this patch is still

>> necessary for case 1.

>>

>> Is this OK for trunk?

>>

>> Cesar

>>

>>

>> trunk-pr78027.diff

>>

>>

>> 2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

>>

>>     PR fortran/78027

>>

>>     gcc/

>>     * ipa-icf.c (sem_function::parse): Don't process functions with

>>     oacc decl attributes, as they may be OpenACC routines.

>>

>>     gcc/testsuite/

>>     * c-c++-common/goacc/acc-icf.c: New test.

>>     * gfortran.dg/goacc/pr78027.f90: New test.

> This follows the same approach as we do for openmp.  This is fine for

> the trunk.


Is this also OK for gcc-6?

Thanks,
Cesar
Alexander Monakov Dec. 9, 2016, 2:56 p.m. UTC | #4
On Thu, 8 Dec 2016, Jeff Law wrote:
> > 	PR fortran/78027

> > 

> > 	gcc/

> > 	* ipa-icf.c (sem_function::parse): Don't process functions with

> > 	oacc decl attributes, as they may be OpenACC routines.

> > 

> > 	gcc/testsuite/

> > 	* c-c++-common/goacc/acc-icf.c: New test.

> > 	* gfortran.dg/goacc/pr78027.f90: New test.

> This follows the same approach as we do for openmp.  This is fine for the

> trunk.


Normally all offloaded code has either "omp declare target" (most functions) or
"omp target entrypoint" (only toplevel offloaded code).  This used to include
OpenACC functions: they would have "oacc fnattrib" _in_ _addition_ to those
attributes, not _instead_ of them!  If something related to a recent OpenACC
patch adds "oacc fnattrib" without at the same time adding "omp declare target",
it's probable that other places in the compiler may break, not just IPA ICF.

The explanation and the patch are written as if "oacc fnattrib" can appear
without "omp declare target" -- ttbomk this could never happen before.

Can this please be cleared up?

Alexander

P.S. (and for the record, I'm unhappy that the reason for blacklisting "omp *"
attributes doesn't seem to be documented anywhere -- I see no way to understand
from looking at this code _why_ the check is there; is it due to issues with
"omp declare simd"? or something else?)
Cesar Philippidis Dec. 9, 2016, 3:37 p.m. UTC | #5
On 12/09/2016 06:56 AM, Alexander Monakov wrote:
> On Thu, 8 Dec 2016, Jeff Law wrote:

>>> 	PR fortran/78027

>>>

>>> 	gcc/

>>> 	* ipa-icf.c (sem_function::parse): Don't process functions with

>>> 	oacc decl attributes, as they may be OpenACC routines.

>>>

>>> 	gcc/testsuite/

>>> 	* c-c++-common/goacc/acc-icf.c: New test.

>>> 	* gfortran.dg/goacc/pr78027.f90: New test.

>> This follows the same approach as we do for openmp.  This is fine for the

>> trunk.

> 

> Normally all offloaded code has either "omp declare target" (most functions) or

> "omp target entrypoint" (only toplevel offloaded code).  This used to include

> OpenACC functions: they would have "oacc fnattrib" _in_ _addition_ to those

> attributes, not _instead_ of them!  If something related to a recent OpenACC

> patch adds "oacc fnattrib" without at the same time adding "omp declare target",

> it's probable that other places in the compiler may break, not just IPA ICF.

> 

> The explanation and the patch are written as if "oacc fnattrib" can appear

> without "omp declare target" -- ttbomk this could never happen before.

> 

> Can this please be cleared up?


It was more so to point out that there were cases where OpenACC
offloaded regions may not have omp target attributes. Referring back to
case 2 in the original email:

  2. Nested offloaded regions inside 'omp declare target' functions do
  not have 'omp target entrypoint' attributes either.

See the logic in omp-low.c:create_omp_child_function. Therefore,
something like pr78027.f90 will not work. As I mentioned earlier, this
probably should be a parser error. Looking back at the spec, it looks
like omp declare target functions cannot have any nested omp target
regions in them. A natural extension would be to error on acc parallel
and kernels regions inside omp declare and omp target regions inside acc
routines.

Neither the OpenMP or OpenACC spec mentions anything about
interoperability with one another. The OpenACC committee is planning on
addressing OpenMP interoperability in 4.0 though. What's a better
solution? A) always add an 'omp target entrypoint' attribute for OpenACC
parallel and kernels regions, B) generate a parser error for
OpenACC/OpenMP target/parallel/kernels regions inside routines/omp
declared functions, C) error on -fopenmp -fopenacc or D) leave this
patch as-is?

> P.S. (and for the record, I'm unhappy that the reason for blacklisting "omp *"

> attributes doesn't seem to be documented anywhere -- I see no way to understand

> from looking at this code _why_ the check is there; is it due to issues with

> "omp declare simd"? or something else?)


That does seem a little strict. I was debating whether this would be
valid for OpenACC or not. The tricky part with OpenACC is the launch
geometry. But then again, we could always teach IPA-ICF how to check the
function attributes for equality among ACC regions.

Cesar
Jakub Jelinek Dec. 9, 2016, 3:48 p.m. UTC | #6
On Fri, Dec 09, 2016 at 07:37:53AM -0800, Cesar Philippidis wrote:
> It was more so to point out that there were cases where OpenACC

> offloaded regions may not have omp target attributes. Referring back to

> case 2 in the original email:

> 

>   2. Nested offloaded regions inside 'omp declare target' functions do

>   not have 'omp target entrypoint' attributes either.

> 

> See the logic in omp-low.c:create_omp_child_function. Therefore,

> something like pr78027.f90 will not work. As I mentioned earlier, this

> probably should be a parser error. Looking back at the spec, it looks

> like omp declare target functions cannot have any nested omp target

> regions in them. A natural extension would be to error on acc parallel

> and kernels regions inside omp declare and omp target regions inside acc

> routines.


#pragma omp target is UB inside of target region, so it is e.g. fine to
turn it into __builtin_trap.  But, rejecting it is problematic.
Warning about it when lexically nested in another #pragma omp target
is fine, when nested in #pragma omp declare target is wrong:
Having a #pragma omp declare target function which contains
#pragma omp target can be valid, as long as you never call the function in
the target region, or as long as when executing that function in the target
region you never invoke the #pragma omp target in there.

	Jakub
Alexander Monakov Dec. 9, 2016, 3:49 p.m. UTC | #7
On Fri, 9 Dec 2016, Cesar Philippidis wrote:
> > Normally all offloaded code has either "omp declare target" (most functions) or

> > "omp target entrypoint" (only toplevel offloaded code).  This used to include

> > OpenACC functions: they would have "oacc fnattrib" _in_ _addition_ to those

> > attributes, not _instead_ of them!  If something related to a recent OpenACC

> > patch adds "oacc fnattrib" without at the same time adding "omp declare target",

> > it's probable that other places in the compiler may break, not just IPA ICF.

> > 

> > The explanation and the patch are written as if "oacc fnattrib" can appear

> > without "omp declare target" -- ttbomk this could never happen before.

> > 

> > Can this please be cleared up?

> 

> It was more so to point out that there were cases where OpenACC

> offloaded regions may not have omp target attributes. Referring back to

> case 2 in the original email:

> 

>   2. Nested offloaded regions inside 'omp declare target' functions do

>   not have 'omp target entrypoint' attributes either.

> 

> See the logic in omp-low.c:create_omp_child_function.


Let's see:

>  if (cgraph_node::get_create (decl)->offloadable

>      && !lookup_attribute ("omp declare target",

>                           DECL_ATTRIBUTES (current_function_decl)))

>    {

>      const char *target_attr = (is_gimple_omp_offloaded (ctx->stmt)

>                                 ? "omp target entrypoint"

>                                 : "omp declare target");

>      DECL_ATTRIBUTES (decl)

>        = tree_cons (get_identifier (target_attr),

>                     NULL_TREE, DECL_ATTRIBUTES (decl));

>    }


So, like I said -- offloaded functions have either "omp target entrypoint" or
"omp declare function".  Either of those are caught by the _existing_ "omp "
prefix check in IPA ICF.

It doesn't matter that nested entrypoints might not have "omp target entrypoint"
on them, because even if they have just "omp declare target" it's still subject
to that ICF blacklist.

My objection still stands -- if you have "oacc fnattr" functions that have
neither of those attributes, that doesn't seem right.

Alexander
Cesar Philippidis Dec. 9, 2016, 4:12 p.m. UTC | #8
On 12/09/2016 07:49 AM, Alexander Monakov wrote:
> On Fri, 9 Dec 2016, Cesar Philippidis wrote:

>>> Normally all offloaded code has either "omp declare target" (most functions) or

>>> "omp target entrypoint" (only toplevel offloaded code).  This used to include

>>> OpenACC functions: they would have "oacc fnattrib" _in_ _addition_ to those

>>> attributes, not _instead_ of them!  If something related to a recent OpenACC

>>> patch adds "oacc fnattrib" without at the same time adding "omp declare target",

>>> it's probable that other places in the compiler may break, not just IPA ICF.

>>>

>>> The explanation and the patch are written as if "oacc fnattrib" can appear

>>> without "omp declare target" -- ttbomk this could never happen before.

>>>

>>> Can this please be cleared up?

>>

>> It was more so to point out that there were cases where OpenACC

>> offloaded regions may not have omp target attributes. Referring back to

>> case 2 in the original email:

>>

>>   2. Nested offloaded regions inside 'omp declare target' functions do

>>   not have 'omp target entrypoint' attributes either.

>>

>> See the logic in omp-low.c:create_omp_child_function.

> 

> Let's see:

> 

>>  if (cgraph_node::get_create (decl)->offloadable

>>      && !lookup_attribute ("omp declare target",

>>                           DECL_ATTRIBUTES (current_function_decl)))

>>    {

>>      const char *target_attr = (is_gimple_omp_offloaded (ctx->stmt)

>>                                 ? "omp target entrypoint"

>>                                 : "omp declare target");

>>      DECL_ATTRIBUTES (decl)

>>        = tree_cons (get_identifier (target_attr),

>>                     NULL_TREE, DECL_ATTRIBUTES (decl));

>>    }

> 

> So, like I said -- offloaded functions have either "omp target entrypoint" or

> "omp declare function".  Either of those are caught by the _existing_ "omp "

> prefix check in IPA ICF.


How did you interpret this from that code?

> It doesn't matter that nested entrypoints might not have "omp target entrypoint"

> on them, because even if they have just "omp declare target" it's still subject

> to that ICF blacklist.

> 

> My objection still stands -- if you have "oacc fnattr" functions that have

> neither of those attributes, that doesn't seem right.


I'm confused. Are you implying that create_omp_child_function is
correct, or are you arguing for omp function attributes to always be
present on all offloaded functions, or both? Looking at this example again:

real function f()
   !$omp declare target(f)
   f = 1.
   !$acc parallel
   !$acc loop
   do i = 1, 8
   end do
   !$acc end parallel
   !$acc parallel
   !$acc loop
   do i = 1, 8
   end do
   !$acc end parallel
 end

Here an 'acc parallel' region is nested inside a function with an 'omp
declare target' attribute. Therefore, because current_function_decl, f,
already has an 'omp target declare' attribute, neither of the acc
parallel regions will get marked with an omp function attribute. Hence
the four solutions I listed earlier.

Cesar
Alexander Monakov Dec. 9, 2016, 5:28 p.m. UTC | #9
On Fri, 9 Dec 2016, Cesar Philippidis wrote:
> >>  if (cgraph_node::get_create (decl)->offloadable

> >>      && !lookup_attribute ("omp declare target",

> >>                           DECL_ATTRIBUTES (current_function_decl)))

> >>    {

> >>      const char *target_attr = (is_gimple_omp_offloaded (ctx->stmt)

> >>                                 ? "omp target entrypoint"

> >>                                 : "omp declare target");

> >>      DECL_ATTRIBUTES (decl)

> >>        = tree_cons (get_identifier (target_attr),

> >>                     NULL_TREE, DECL_ATTRIBUTES (decl));

> >>    }

> > 

> > So, like I said -- offloaded functions have either "omp target entrypoint" or

> > "omp declare function".  Either of those are caught by the _existing_ "omp "

> > prefix check in IPA ICF.

> 

> How did you interpret this from that code?


Bah, sorry, I was missing the distinction between 'decl' and
'current_function_decl' and was always interpreting this code like 'add an
attribute, unless "omp declare function" is already present on decl'.

Perhaps the original intention was to suppress creating target entrypoints
inside of offloaded code.  But in that case, the code should have checked either
"omp declare target" or "omp target entrypoint" on the containing function.  I
think this is quite confusing.  I've changed this code in OpenMP/NVPTX work, but
not too much.

Jakub -- do you think I'm deducing the original intent here correctly?  Can we
a) fix that code to properly follow that intention, and
b) start enforcing the invariant that all offloaded code has either "omp target
entrypoint" or "omp declare function"?

> > on them, because even if they have just "omp declare target" it's still subject

> > to that ICF blacklist.

> > 

> > My objection still stands -- if you have "oacc fnattr" functions that have

> > neither of those attributes, that doesn't seem right.

> 

> I'm confused. Are you implying that create_omp_child_function is

> correct, or are you arguing for omp function attributes to always be

> present on all offloaded functions, or both?


The 2nd option -- I thought offloaded functions were intended to have one of
"omp *" attributes, but it appears I might have been mistaken.


Still, "offloaded functions must have one of two omp attributes" sounds like a
simpler invariant than "<the above> or the oacc attribute [and all code checking
the former must start checking the oacc attribute too]".

Thanks for bearing with me.
Alexander
Jeff Law Dec. 12, 2016, 4:29 p.m. UTC | #10
On 12/09/2016 07:05 AM, Cesar Philippidis wrote:
> On 12/08/2016 04:37 PM, Jeff Law wrote:

>> On 12/08/2016 04:05 PM, Cesar Philippidis wrote:

>>> PR78027 was classified as a fortran bug, but the underlying problem

>>> turned out to be more generic. Basically, the IPA-ICF pass usually

>>> ignores functions with omp decl attributes as candidates for aliasing.

>>> Usually that works for OpenACC, because offloaded functions generated by

>>> OpenACC PARALLEL or KERNELS regions usually have an 'omp target

>>> entrypoint' attribute. However that is not the case in two situations:

>>>

>>> 1. ACC ROUTINES do not have an 'omp target entrypoint' attribute.

>>>

>>> 2. Nested offloaded regions inside 'omp declare target' functions do not

>>> have 'omp target entrypoint' attributes either.

>>>

>>> The solution I chose for this problem is to teach the IPA-ICF pass to

>>> ignore function with 'oacc *' decl attributes. Arguably 2) should be a

>>> parser error when an OpenACC offloaded region is embedded within an

>>> OpenMP offloaded function. The LTO linker does report an error for nvptx

>>> targets, but not the host. With that in mind, this patch is still

>>> necessary for case 1.

>>>

>>> Is this OK for trunk?

>>>

>>> Cesar

>>>

>>>

>>> trunk-pr78027.diff

>>>

>>>

>>> 2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

>>>

>>>     PR fortran/78027

>>>

>>>     gcc/

>>>     * ipa-icf.c (sem_function::parse): Don't process functions with

>>>     oacc decl attributes, as they may be OpenACC routines.

>>>

>>>     gcc/testsuite/

>>>     * c-c++-common/goacc/acc-icf.c: New test.

>>>     * gfortran.dg/goacc/pr78027.f90: New test.

>> This follows the same approach as we do for openmp.  This is fine for

>> the trunk.

>

> Is this also OK for gcc-6?

Your call -- I don't think folks are really using OpenACC much in gcc-6, 
but if you think it's worth the time, go ahead.

jeff
Jakub Jelinek Dec. 12, 2016, 4:37 p.m. UTC | #11
On Mon, Dec 12, 2016 at 09:29:44AM -0700, Jeff Law wrote:
> >>>2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

> >>>

> >>>    PR fortran/78027

> >>>

> >>>    gcc/

> >>>    * ipa-icf.c (sem_function::parse): Don't process functions with

> >>>    oacc decl attributes, as they may be OpenACC routines.

> >>>

> >>>    gcc/testsuite/

> >>>    * c-c++-common/goacc/acc-icf.c: New test.

> >>>    * gfortran.dg/goacc/pr78027.f90: New test.

> >>This follows the same approach as we do for openmp.  This is fine for

> >>the trunk.

> >

> >Is this also OK for gcc-6?

> Your call -- I don't think folks are really using OpenACC much in gcc-6, but

> if you think it's worth the time, go ahead.


As Alex noted, the patch looks wrong, there should be "omp declare target"
attribute, which isn't specific to OpenMP only, but covers all the
offloading fns.  I have on my todo list to look at it.

	Jakub
Jeff Law Dec. 12, 2016, 4:38 p.m. UTC | #12
On 12/12/2016 09:37 AM, Jakub Jelinek wrote:
> On Mon, Dec 12, 2016 at 09:29:44AM -0700, Jeff Law wrote:

>>>>> 2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

>>>>>

>>>>>    PR fortran/78027

>>>>>

>>>>>    gcc/

>>>>>    * ipa-icf.c (sem_function::parse): Don't process functions with

>>>>>    oacc decl attributes, as they may be OpenACC routines.

>>>>>

>>>>>    gcc/testsuite/

>>>>>    * c-c++-common/goacc/acc-icf.c: New test.

>>>>>    * gfortran.dg/goacc/pr78027.f90: New test.

>>>> This follows the same approach as we do for openmp.  This is fine for

>>>> the trunk.

>>>

>>> Is this also OK for gcc-6?

>> Your call -- I don't think folks are really using OpenACC much in gcc-6, but

>> if you think it's worth the time, go ahead.

>

> As Alex noted, the patch looks wrong, there should be "omp declare target"

> attribute, which isn't specific to OpenMP only, but covers all the

> offloading fns.  I have on my todo list to look at it.

I thought the final resolution of that thread was a misunderstanding 
from Alex as to how this stuff worked.  If I'm wrong, then obviously we 
shouldn't apply the patch to the trunk or branch.

jeff
Thomas Schwinge Feb. 22, 2019, 10:58 a.m. UTC | #13
Hi!

On Thu, 8 Dec 2016 15:05:10 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- /dev/null

> +++ b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90

> @@ -0,0 +1,20 @@

> +! { dg-additional-options "-fopenmp -O2 -fdump-ipa-icf" }

> +[...]


To address unrelated compiler diagnostics, I've just committed to trunk
in r269103 "Silence '-Whsa' diagnostic in
'gfortran.dg/goacc/pr78027.f90'", as attached.


Grüße
 Thomas
From eb1dc7b8b939858939783d0ce5ba6d5375784dd1 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>

Date: Fri, 22 Feb 2019 10:50:02 +0000
Subject: [PATCH 2/9] Silence '-Whsa' diagnostic in
 'gfortran.dg/goacc/pr78027.f90'

... which has been present (with HSA offloading configured) ever since this
test case got added.

	gcc/testsuite/
	PR fortran/78027
	* gfortran.dg/goacc/pr78027.f90: Add 'dg-additional-options "-Wno-hsa"'.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@269103 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                     | 5 +++++
 gcc/testsuite/gfortran.dg/goacc/pr78027.f90 | 4 ++++
 2 files changed, 9 insertions(+)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 7cb6b240d0c..b98cd355645 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2019-02-22  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR fortran/78027
+	* gfortran.dg/goacc/pr78027.f90: Add 'dg-additional-options "-Wno-hsa"'.
+
 2019-02-22  Richard Biener  <rguenther@suse.de>
 
 	PR middle-end/87609
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78027.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90
index db65063bede..a1e51fa0f0d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr78027.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90
@@ -1,5 +1,9 @@
 ! { dg-additional-options "-fopenmp -O2 -fdump-ipa-icf" }
 
+!     f951: warning: could not emit HSAIL for the function [-Whsa]
+!     f951: note: HSA does not support functions with variadic arguments (or unknown return type): 'GOACC_parallel_keyed'
+! { dg-additional-options "-Wno-hsa" }
+
 real function f()
    !$omp declare target(f)
    f = 1.
-- 
2.17.1
diff mbox

Patch

2016-12-08  Cesar Philippidis  <cesar@codesourcery.com>

	PR fortran/78027

	gcc/
	* ipa-icf.c (sem_function::parse): Don't process functions with
	oacc decl attributes, as they may be OpenACC routines.

	gcc/testsuite/
	* c-c++-common/goacc/acc-icf.c: New test.
	* gfortran.dg/goacc/pr78027.f90: New test.


diff --git a/gcc/ipa-icf.c b/gcc/ipa-icf.c
index 553b81e..31061e9 100644
--- a/gcc/ipa-icf.c
+++ b/gcc/ipa-icf.c
@@ -1689,6 +1689,10 @@  sem_function::parse (cgraph_node *node, bitmap_obstack *stack)
   if (lookup_attribute_by_prefix ("omp ", DECL_ATTRIBUTES (node->decl)) != NULL)
     return NULL;
 
+  if (lookup_attribute_by_prefix ("oacc ",
+				  DECL_ATTRIBUTES (node->decl)) != NULL)
+    return NULL;
+
   /* PR ipa/70306.  */
   if (DECL_STATIC_CONSTRUCTOR (node->decl)
       || DECL_STATIC_DESTRUCTOR (node->decl))
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
new file mode 100644
index 0000000..ecfe3f2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c
@@ -0,0 +1,49 @@ 
+/* Ensure that IPA-ICF is disabled on OpenACC routines.  */
+
+/* { dg-additional-options "-fopenacc -O2 -fdump-ipa-icf" }  */
+
+#pragma acc routine gang
+int
+routine1 (int n)
+{
+  int i;
+
+  #pragma acc loop
+  for (i = 0; i < n; i++)
+    ;
+
+  return n + 1;
+}
+
+#pragma acc routine gang
+int
+routine2 (int n)
+{
+  int i;
+
+  #pragma acc loop
+  for (i = 0; i < n; i++)
+    ;
+
+  return n + 1;
+}
+
+int
+main ()
+{
+  int i;
+
+  #pragma acc parallel loop
+  for (i = 0; i < 8; i++)
+    ;
+
+  #pragma acc parallel loop
+  for (i = 0; i < 8; i++)
+    ;
+
+  return 0;
+}
+
+/* { dg-final { scan-ipa-dump-times "Not parsed function:" 4 "icf" } }  */
+/* { dg-final { scan-ipa-dump "Parsed function:main" "icf" } }  */
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78027.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90
new file mode 100644
index 0000000..db65063
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90
@@ -0,0 +1,20 @@ 
+! { dg-additional-options "-fopenmp -O2 -fdump-ipa-icf" }
+
+real function f()
+   !$omp declare target(f)
+   f = 1.
+   !$acc parallel
+   !$acc loop
+   do i = 1, 8
+   end do
+   !$acc end parallel
+   !$acc parallel
+   !$acc loop
+   do i = 1, 8
+   end do
+   !$acc end parallel
+ end
+ 
+! { dg-final { scan-ipa-dump "Not parsed function:f_._omp_fn.1" "icf" } }
+! { dg-final { scan-ipa-dump "Not parsed function:f_._omp_fn.0" "icf" } }
+! { dg-final { scan-ipa-dump "Not parsed function:f_" "icf" } }