Hi Jakub,
Jakub Jelinek wrote:
diff --git a/libgomp/testsuite/libgomp.c++/declare_target-2.C
b/libgomp/testsuite/libgomp.c++/declare_target-2.C
+// Actually not needed: -fipa-cp is default with -O2:
+// { dg-additional-options "-O2 -fipa-cp -foffload=amdgcn-amdhsa" }
But am not sure how this will work if amdgcn offloading is not configured.
That's because I messed up during testing to simulate an only-amdgcn
offload. The test is intended to work everywhere, trivially without
offloading but also with nvptx and/or amdgcn offloading enabled.
It should have been removed before doing another 'git add -u' for other
fixes :-/
+// Nvptx fails because of PR libstdc++/120680
+// { dg-bogus "_ZN9__gnu_cxx21zoneinfo_dir_overrideEv" "" { xfail
offload_target_nvptx } 0 }
+// { dg-xfail-if "PR120680" { offload_target_nvptx } }
And why are the xfails needed for nvptx when -foffload is passed for some
other offload target.
The plan (and as tested), it should also work with nvptx offloading
(alone or mixed with gcn) without ICE, but it failed here at link time.
However, it turned out that I only ran into PR120680 because I used a
too-old distro's nvptx-tools, lackting sufficient 'weak' support. With
nvptx-tools > early August 2025, it links. (Too old in 2 of the 3 major
distros; in one, it is new enough, but also misses the Sept 2024 commits.)
Thanks for proof reading! I have now committed the attached patch, which
assumes a sufficiently new nvptx-tools (and has no xfails).
→ r16-1536-gea43b99537591b
Tobias
commit ea43b99537591b1103da3961c61f1cbfae968859
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.
---
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 da2b54b7648..d88edb29a0e 100644
--- 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
index 00000000000..ab94a556888
--- /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);
+}
