OpenMP: Fix implicit 'declare target' for <ostream>

libstdc++-v3/include/std/ostream contains:

  namespace std _GLIBCXX_VISIBILITY(default)
  {
    ...
    template<typename _CharT, typename _Traits>
      inline basic_ostream<_CharT, _Traits>&
      endl(basic_ostream<_CharT, _Traits>& __os)
      { return flush(__os.put(__os.widen('\n'))); }
  ...
  #include <bits/ostream.tcc>

and the latter, libstdc++-v3/include/bits/ostream.tcc, has:
    // Inhibit implicit instantiations for required instantiations,
    // which are defined via explicit instantiations elsewhere.
  #if _GLIBCXX_EXTERN_TEMPLATE
    extern template class basic_ostream<char>;
    extern template ostream& endl(ostream&);

Before this commit, omp_discover_declare_target_tgt_fn_r marked 'endl'
as (implicitly) declare target - but not the calls in it due to the
'extern' (DECL_EXTERNAL).

Thanks to inlining and as 'endl' is (therefore) not used and, hence,
discarded by the linker; hencet, it works with -O0 and -O1. However,
as the (unused) function still exits, IPA CP (enabled by -O2) will try
to do constant-value propagation and fails as the definition of 'widen'
is not available.

Solution is to still walk 'endl' despite being an 'extern(al)' decl;
this has been restricted for now to DECL_DECLARED_INLINE_P.

gcc/ChangeLog:

	* omp-offload.cc (omp_discover_declare_target_tgt_fn_r): Also
	walk external functions that are declare inline (and have a
	DECL_SAVED_TREE).

libgomp/ChangeLog:

	* testsuite/libgomp.c++/declare_target-2.C: New test.

diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index da2b54b764856af649bdab642c6e1367519fb81a..d88edb29a0e83be8cc1b67c08d648ffdf18d1003 100644 (file)
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -261,7 +261,8 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
                               DECL_ATTRIBUTES (decl)))
        return NULL_TREE;
 
-      if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl))
+      if (DECL_SAVED_TREE (decl)
+         && (!DECL_EXTERNAL (decl) || DECL_DECLARED_INLINE_P (decl)))
        ((vec<tree> *) data)->safe_push (decl);
       DECL_ATTRIBUTES (decl) = tree_cons (id, NULL_TREE,
                                          DECL_ATTRIBUTES (decl));

diff --git a/libgomp/testsuite/libgomp.c++/declare_target-2.C b/libgomp/testsuite/libgomp.c++/declare_target-2.C
new file mode 100644 (file)
index 0000000..ab94a55
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare_target-2.C
@@ -0,0 +1,25 @@
+// { dg-do link }
+
+// Actually not needed: -fipa-cp is default with -O2:
+// { dg-additional-options "-O2 -fipa-cp" }
+
+// The code failed because 'std::endl' becoḿes implicitly 'declare target'
+// but not the 'widen' function it calls.  While the linker had no issues
+// (endl is never called, either because it is inlined or optimized away),
+// the IPA-CP (enabled by -O2 and higher) failed as the definition for
+// 'widen' did not exist on the offload side.
+
+#include <iostream>
+
+void func (int m)
+{
+  if (m < 0)
+    std::cout << "should not happen" << std::endl;
+}
+
+
+int main()
+{
+  #pragma omp target
+    func (1);
+}