This fixes an issue with implicit declare target.
'<ostream>' has:
template … endl() { ... }
...
extern template ostream& endl(ostream&);
which leads to a decl that is 'extern' but still
has the full function body available.
Result: On the offload side, 'endl' is present
but the functions it calls are not. Once IPA CP
kicks in, this causes an ICE as called function is
not available, resulting an ICE as some node info
is NULL but dereferenced.
I am not quite sure when a function body should
be walked for implicit declare target, when it
has the DECL_EXTERNAL but still a DECL_SAVED_TREE:
Always or - as with this patch - only if it has the
DECL_DECLARED_INLINE_P property?
Comments, suggestions?
Tobias
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 | 30 ++++++++++++++++++++++++
2 files changed, 32 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..6d6c94c514d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare_target-2.C
@@ -0,0 +1,30 @@
+// { dg-do link }
+
+// Actually not needed: -fipa-cp is default with -O2:
+// { dg-additional-options "-O2 -fipa-cp -foffload=amdgcn-amdhsa" }
+
+// 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.
+
+// 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 } }
+
+
+#include <iostream>
+
+void func (int m)
+{
+ if (m < 0)
+ std::cout << "should not happen" << std::endl;
+}
+
+
+int main()
+{
+ #pragma omp target
+ func (1);
+}
