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
2016-12-08 Cesar Philippidis <ce...@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" } }
