[gcc/devel/omp/gcc-15] OpenMP: Fix implicit 'declare target' for

[Tobias Burnus via Gcc-cvs]

https://gcc.gnu.org/g:9ee74468c5a145b0926609d3066af059eeb2ad0c

commit 9ee74468c5a145b0926609d3066af059eeb2ad0c
Author: Tobias Burnus <tbur...@baylibre.com>
Date:   Tue Jun 17 11:33:09 2025 +0200

    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)

Diff:
---
 gcc/omp-offload.cc                               |  3 ++-
 libgomp/testsuite/libgomp.c++/declare_target-2.C | 25 ++++++++++++++++++++++++
 2 files changed, 27 insertions(+), 1 deletion(-)

diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 3218f6976f3f..16fa232352f9 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -262,7 +262,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
index 000000000000..ab94a5568881
--- /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);
+}