commit 9023bcd0d68c1ba7faebc9edb6db34537b2a7966
author Tobias Burnus <tburnus@baylibre.com> Tue Jun 17 11:33:09 2025 +0200
committer Thomas Schwinge <tschwinge@baylibre.com> Tue Jul 22 11:11:51 2025 +0200
tree dc1d7db1540d13c5e5df76452107f68f6ceb962d
parent dd10f193a877a45c16682b4b9ee58bdc281859ff

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.

(cherry picked from commit ea43b99537591b1103da3961c61f1cbfae968859)

gcc/omp-offload.cc

diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index da2b54b..d88edb2 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -261,7 +261,8 @@
 			       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));

libgomp/testsuite/libgomp.c++/declare_target-2.C

diff --git a/libgomp/testsuite/libgomp.c++/declare_target-2.C b/libgomp/testsuite/libgomp.c++/declare_target-2.C
new file mode 100644
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);
+}