https://gcc.gnu.org/g:fe283dba774be57b705a7a871b000d2894d2e553
commit r15-9470-gfe283dba774be57b705a7a871b000d2894d2e553
Author: Thomas Schwinge <tschwi...@baylibre.com>
Date: Fri Mar 28 09:20:49 2025 +0100
GCN, nvptx: Support '-mfake-exceptions', and use it for offloading
compilation [PR118794]
With '-mfake-exceptions' enabled, the user-visible behavior in presence of
exception handling constructs changes such that the compile-time
'sorry, unimplemented: exception handling not supported' is skipped, code
generation proceeds, and instead, exception handling constructs 'abort' at
run time. (..., or don't, if they're in dead code.)
PR target/118794
gcc/
* config/gcn/gcn.opt (-mfake-exceptions): Support.
* config/nvptx/nvptx.opt (-mfake-exceptions): Likewise.
* config/gcn/gcn.md (define_expand "exception_receiver"): Use it.
* config/nvptx/nvptx.md (define_expand "exception_receiver"):
Likewise.
* config/gcn/mkoffload.cc (main): Set it.
* config/nvptx/mkoffload.cc (main): Likewise.
* config/nvptx/nvptx.cc (nvptx_assemble_integer)
<in_section == exception_section>: Special handling for
'SYMBOL_REF's.
* except.cc (expand_dw2_landing_pad_for_region): Don't generate
bogus code for (default)
'#define EH_RETURN_DATA_REGNO(N) INVALID_REGNUM'.
libgcc/
* config/gcn/unwind-gcn.c (_Unwind_Resume): New.
* config/nvptx/unwind-nvptx.c (_Unwind_Resume): Likewise.
gcc/testsuite/
* g++.target/gcn/exceptions-bad_cast-2.C: Set
'-mno-fake-exceptions'.
* g++.target/gcn/exceptions-pr118794-1.C: Likewise.
* g++.target/gcn/exceptions-throw-2.C: Likewise.
* g++.target/nvptx/exceptions-bad_cast-2.C: Likewise.
* g++.target/nvptx/exceptions-pr118794-1.C: Likewise.
* g++.target/nvptx/exceptions-throw-2.C: Likewise.
* g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C: New.
* g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C:
Likewise.
* g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C: Likewise.
* g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C:
Likewise.
* g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C:
Likewise.
* g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C:
Likewise.
libgomp/
*
testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C:
Set '-foffload-options=-mno-fake-exceptions'.
*
testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C:
Likewise.
*
testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C:
Likewise.
*
testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C:
Likewise.
*
testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C:
Likewise.
*
testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C:
Likewise.
*
testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C:
Likewise.
*
testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C:
Likewise.
* testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C:
Likewise.
*
testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C:
Likewise.
* testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Adjust.
* testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-throw-2.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-throw-2-O0.C: New.
Diff:
---
gcc/config/gcn/gcn.md | 4 ++-
gcc/config/gcn/gcn.opt | 8 ++++++
gcc/config/gcn/mkoffload.cc | 3 +++
gcc/config/nvptx/mkoffload.cc | 3 +++
gcc/config/nvptx/nvptx.cc | 20 ++++++++++++++-
gcc/config/nvptx/nvptx.md | 4 ++-
gcc/config/nvptx/nvptx.opt | 8 ++++++
gcc/except.cc | 24 +++++++++++++----
.../g++.target/gcn/exceptions-bad_cast-2.C | 1 +
.../gcn/exceptions-bad_cast-2_-mfake-exceptions.C | 18 +++++++++++++
.../g++.target/gcn/exceptions-pr118794-1.C | 1 +
.../gcn/exceptions-pr118794-1_-mfake-exceptions.C | 16 ++++++++++++
gcc/testsuite/g++.target/gcn/exceptions-throw-2.C | 1 +
.../gcn/exceptions-throw-2_-mfake-exceptions.C | 19 ++++++++++++++
.../g++.target/nvptx/exceptions-bad_cast-2.C | 1 +
.../exceptions-bad_cast-2_-mfake-exceptions.C | 19 ++++++++++++++
.../g++.target/nvptx/exceptions-pr118794-1.C | 1 +
.../exceptions-pr118794-1_-mfake-exceptions.C | 16 ++++++++++++
.../g++.target/nvptx/exceptions-throw-2.C | 1 +
.../nvptx/exceptions-throw-2_-mfake-exceptions.C | 19 ++++++++++++++
libgcc/config/gcn/unwind-gcn.c | 6 +++++
libgcc/config/nvptx/unwind-nvptx.c | 6 +++++
...arget-exceptions-bad_cast-2-offload-sorry-GCN.C | 5 ++--
...get-exceptions-bad_cast-2-offload-sorry-nvptx.C | 5 ++--
.../libgomp.c++/target-exceptions-bad_cast-2.C | 24 ++++++++++++-----
...arget-exceptions-pr118794-1-offload-sorry-GCN.C | 5 ++--
...get-exceptions-pr118794-1-offload-sorry-nvptx.C | 5 ++--
.../libgomp.c++/target-exceptions-pr118794-1.C | 17 +++++++-----
.../libgomp.c++/target-exceptions-throw-2-O0.C | 25 ++++++++++++++++++
.../target-exceptions-throw-2-offload-sorry-GCN.C | 5 ++--
...target-exceptions-throw-2-offload-sorry-nvptx.C | 5 ++--
.../libgomp.c++/target-exceptions-throw-2.C | 22 +++++++++++-----
.../exceptions-bad_cast-2-offload-sorry-GCN.C | 5 ++--
.../exceptions-bad_cast-2-offload-sorry-nvptx.C | 5 ++--
.../libgomp.oacc-c++/exceptions-bad_cast-2.C | 26 ++++++++++++++-----
.../exceptions-throw-2-offload-sorry-GCN.C | 5 ++--
.../exceptions-throw-2-offload-sorry-nvptx.C | 5 ++--
.../libgomp.oacc-c++/exceptions-throw-2.C | 30 +++++++++++++++++-----
38 files changed, 331 insertions(+), 62 deletions(-)
diff --git a/gcc/config/gcn/gcn.md b/gcc/config/gcn/gcn.md
index 695656f692d6..e0fb735e5ec9 100644
--- a/gcc/config/gcn/gcn.md
+++ b/gcc/config/gcn/gcn.md
@@ -1018,7 +1018,9 @@
[(const_int 0)]
""
{
- sorry ("exception handling not supported");
+ if (!fake_exceptions)
+ sorry ("exception handling not supported");
+ DONE;
})
;; }}}
diff --git a/gcc/config/gcn/gcn.opt b/gcc/config/gcn/gcn.opt
index 142b439bbb88..99d6aeb2b30f 100644
--- a/gcc/config/gcn/gcn.opt
+++ b/gcc/config/gcn/gcn.opt
@@ -101,3 +101,11 @@ Enum(gcn_preferred_vectorization_factor) String(32)
Value(32)
EnumValue
Enum(gcn_preferred_vectorization_factor) String(64) Value(64)
+
+mfake-exceptions
+Target Var(fake_exceptions) Init(0) Undocumented
+; With '-mfake-exceptions' enabled, the user-visible behavior in presence of
+; exception handling constructs changes such that the compile-time
+; 'sorry, unimplemented: exception handling not supported' is skipped, code
+; generation proceeds, and instead, exception handling constructs 'abort' at
+; run time. (..., or don't, if they're in dead code.)
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index f5b89c9fee7a..b284ff422ee7 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -1160,6 +1160,9 @@ main (int argc, char **argv)
obstack_ptr_grow (&cc_argv_obstack, "-xlto");
if (fopenmp)
obstack_ptr_grow (&cc_argv_obstack, "-mgomp");
+ /* The host code may contain exception handling constructs.
+ Handle these as good as we can. */
+ obstack_ptr_grow (&cc_argv_obstack, "-mfake-exceptions");
for (int ix = 1; ix != argc; ix++)
{
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index bdfe7f58b166..e7ec0ef4f6ac 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -778,6 +778,9 @@ main (int argc, char **argv)
}
if (fopenmp)
obstack_ptr_grow (&argv_obstack, "-mgomp");
+ /* The host code may contain exception handling constructs.
+ Handle these as good as we can. */
+ obstack_ptr_grow (&argv_obstack, "-mfake-exceptions");
for (int ix = 1; ix != argc; ix++)
{
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index 87364bf423c5..28da43ca740b 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -2359,7 +2359,25 @@ nvptx_assemble_integer (rtx x, unsigned int size, int
ARG_UNUSED (aligned_p))
{
gcc_checking_assert (!init_frag.active);
/* Just use the default machinery; it's not getting used, anyway. */
- return default_assemble_integer (x, size, aligned_p);
+ bool ok = default_assemble_integer (x, size, aligned_p);
+ /* ..., but a few cases need special handling. */
+ switch (GET_CODE (x))
+ {
+ case SYMBOL_REF:
+ /* The default machinery won't work: we don't define the necessary
+ operations; don't use them outside of this. */
+ gcc_checking_assert (!ok);
+ {
+ /* Just emit something; it's not getting used, anyway. */
+ const char *op = "\t.symbol_ref\t";
+ ok = (assemble_integer_with_op (op, x), true);
+ }
+ break;
+
+ default:
+ break;
+ }
+ return ok;
}
gcc_checking_assert (init_frag.active);
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 3201247c5fb1..7c3bd69cc973 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -1644,7 +1644,9 @@
[(const_int 0)]
""
{
- sorry ("exception handling not supported");
+ if (!fake_exceptions)
+ sorry ("exception handling not supported");
+ DONE;
})
(define_expand "nonlocal_goto"
diff --git a/gcc/config/nvptx/nvptx.opt b/gcc/config/nvptx/nvptx.opt
index 9be81aecdddf..ce9fbc7312e0 100644
--- a/gcc/config/nvptx/nvptx.opt
+++ b/gcc/config/nvptx/nvptx.opt
@@ -168,6 +168,14 @@ Target Var(nvptx_alias) Init(0) Undocumented
mexperimental
Target Var(nvptx_experimental) Init(0) Undocumented
+mfake-exceptions
+Target Var(fake_exceptions) Init(0) Undocumented
+; With '-mfake-exceptions' enabled, the user-visible behavior in presence of
+; exception handling constructs changes such that the compile-time
+; 'sorry, unimplemented: exception handling not supported' is skipped, code
+; generation proceeds, and instead, exception handling constructs 'abort' at
+; run time. (..., or don't, if they're in dead code.)
+
mfake-ptx-alloca
Target Var(nvptx_fake_ptx_alloca) Init(0) Undocumented
; With '-mfake-ptx-alloca' enabled, the user-visible behavior changes only
diff --git a/gcc/except.cc b/gcc/except.cc
index d5eb9274a62c..205811c65673 100644
--- a/gcc/except.cc
+++ b/gcc/except.cc
@@ -970,12 +970,26 @@ expand_dw2_landing_pad_for_region (eh_region region)
{ /* Nothing */ }
if (region->exc_ptr_reg)
- emit_move_insn (region->exc_ptr_reg,
- gen_rtx_REG (ptr_mode, EH_RETURN_DATA_REGNO (0)));
+ {
+ rtx exc_ptr_reg;
+ if (EH_RETURN_DATA_REGNO (0) != INVALID_REGNUM)
+ exc_ptr_reg = gen_rtx_REG (ptr_mode, EH_RETURN_DATA_REGNO (0));
+ else
+ /* The target must be doing something special. Submit a dummy. */
+ exc_ptr_reg = constm1_rtx;
+ emit_move_insn (region->exc_ptr_reg, exc_ptr_reg);
+ }
if (region->filter_reg)
- emit_move_insn (region->filter_reg,
- gen_rtx_REG (targetm.eh_return_filter_mode (),
- EH_RETURN_DATA_REGNO (1)));
+ {
+ rtx filter_reg;
+ if (EH_RETURN_DATA_REGNO (1) != INVALID_REGNUM)
+ filter_reg = gen_rtx_REG (targetm.eh_return_filter_mode (),
+ EH_RETURN_DATA_REGNO (1));
+ else
+ /* The target must be doing something special. Submit a dummy. */
+ filter_reg = constm1_rtx;
+ emit_move_insn (region->filter_reg, filter_reg);
+ }
}
/* Expand the extra code needed at landing pads for dwarf2 unwinding. */
diff --git a/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C
b/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C
index ffe9e19a281e..b047cbed444e 100644
--- a/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C
+++ b/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
#include
"../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C"
diff --git
a/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C
b/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C
new file mode 100644
index 000000000000..2904188ddbed
--- /dev/null
+++ b/gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C
@@ -0,0 +1,18 @@
+/* 'std::bad_cast' exception, caught, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "exceptions-bad_cast-2.C"
+
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ We don't print anything, but just 'abort'.
+
+ There is no 'catch'ing; any exception is fatal.
+ { dg-shouldfail {'std::bad_cast' exception} } */
diff --git a/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C
b/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C
index a7013791511d..20f9d49004cd 100644
--- a/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C
+++ b/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -O0 } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
diff --git
a/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C
b/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C
new file mode 100644
index 000000000000..a5f0da2bca7d
--- /dev/null
+++ b/gcc/testsuite/g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C
@@ -0,0 +1,16 @@
+/* Exception handling constructs in dead code, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -O0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "exceptions-pr118794-1.C"
+
+/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes
+ '__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f':
+ { dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1
optimized } }
+ { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } } */
diff --git a/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
b/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
index 0ae64eebd478..671c810a5628 100644
--- a/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
+++ b/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
#include "../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C"
diff --git
a/gcc/testsuite/g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C
b/gcc/testsuite/g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C
new file mode 100644
index 000000000000..f1fd5058cbdf
--- /dev/null
+++ b/gcc/testsuite/g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C
@@ -0,0 +1,19 @@
+/* 'throw', caught, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "exceptions-throw-2.C"
+
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
+ We don't print anything, but just 'abort'.
+
+ There is no 'catch'ing; any exception is fatal.
+ { dg-shouldfail {'MyException' exception} } */
diff --git a/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C
b/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C
index ffe9e19a281e..b047cbed444e 100644
--- a/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
#include
"../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C"
diff --git
a/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C
b/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C
new file mode 100644
index 000000000000..3f40951b278b
--- /dev/null
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C
@@ -0,0 +1,19 @@
+/* 'std::bad_cast' exception, caught, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-bogus {_ZTISt8bad_cast} PR119734 { xfail *-*-* } 0 } */
+
+#include "exceptions-bad_cast-2.C"
+
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ We don't print anything, but just 'abort'.
+
+ There is no 'catch'ing; any exception is fatal.
+ { dg-shouldfail {'std::bad_cast' exception} } */
diff --git a/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C
b/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C
index a7013791511d..20f9d49004cd 100644
--- a/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -O0 } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
diff --git
a/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C
b/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C
new file mode 100644
index 000000000000..a5f0da2bca7d
--- /dev/null
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C
@@ -0,0 +1,16 @@
+/* Exception handling constructs in dead code, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -O0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "exceptions-pr118794-1.C"
+
+/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes
+ '__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f':
+ { dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1
optimized } }
+ { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } } */
diff --git a/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
index 0ae64eebd478..671c810a5628 100644
--- a/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
@@ -2,6 +2,7 @@
/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw } */
#include "../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C"
diff --git
a/gcc/testsuite/g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C
b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C
new file mode 100644
index 000000000000..f1fd5058cbdf
--- /dev/null
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C
@@ -0,0 +1,19 @@
+/* 'throw', caught, '-mfake-exceptions'. */
+
+/* { dg-do run } */
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++
standard is sufficient. */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -mfake-exceptions }
+ { dg-bogus {sorry, unimplemented: exception handling not supported} {} {
target *-*-* } 0 } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "exceptions-throw-2.C"
+
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
+ We don't print anything, but just 'abort'.
+
+ There is no 'catch'ing; any exception is fatal.
+ { dg-shouldfail {'MyException' exception} } */
diff --git a/libgcc/config/gcn/unwind-gcn.c b/libgcc/config/gcn/unwind-gcn.c
index eae741c55bbb..97e22c035b3d 100644
--- a/libgcc/config/gcn/unwind-gcn.c
+++ b/libgcc/config/gcn/unwind-gcn.c
@@ -38,6 +38,12 @@ _Unwind_DeleteException (struct _Unwind_Exception *exc)
(*exc->exception_cleanup) (_URC_FOREIGN_EXCEPTION_CAUGHT, exc);
}
+void
+_Unwind_Resume (struct _Unwind_Exception *exc __attribute__ ((__unused__)))
+{
+ __builtin_abort ();
+}
+
_Unwind_Reason_Code
_Unwind_Resume_or_Rethrow (struct _Unwind_Exception *exc __attribute__
((__unused__)))
{
diff --git a/libgcc/config/nvptx/unwind-nvptx.c
b/libgcc/config/nvptx/unwind-nvptx.c
index eae741c55bbb..97e22c035b3d 100644
--- a/libgcc/config/nvptx/unwind-nvptx.c
+++ b/libgcc/config/nvptx/unwind-nvptx.c
@@ -38,6 +38,12 @@ _Unwind_DeleteException (struct _Unwind_Exception *exc)
(*exc->exception_cleanup) (_URC_FOREIGN_EXCEPTION_CAUGHT, exc);
}
+void
+_Unwind_Resume (struct _Unwind_Exception *exc __attribute__ ((__unused__)))
+{
+ __builtin_abort ();
+}
+
_Unwind_Reason_Code
_Unwind_Resume_or_Rethrow (struct _Unwind_Exception *exc __attribute__
((__unused__)))
{
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C
index ed670a80fb63..93884dfdbf84 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C
@@ -1,10 +1,11 @@
-/* 'std::bad_cast' exception in OpenMP 'target' region, caught. */
+/* 'std::bad_cast' exception in OpenMP 'target' region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_amdgcn } }
{ dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -12,7 +13,7 @@
/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C
index 84506304f121..83ec89bde8c1 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C
@@ -1,10 +1,11 @@
-/* 'std::bad_cast' exception in OpenMP 'target' region, caught. */
+/* 'std::bad_cast' exception in OpenMP 'target' region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_nvptx } }
{ dg-additional-options -foffload=nvptx-none } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -12,7 +13,7 @@
/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C
index 26decf60055d..88617400a723 100644
--- a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C
@@ -2,13 +2,23 @@
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
-/* { dg-additional-options -foffload=disable }
- Offloading compilation not yet supported; see
- 'target-exceptions-bad_cast-2-offload-sorry-GCN.C',
- 'target-exceptions-bad_cast-2-offload-sorry-nvptx.C'. */
-/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+/* { dg-bogus {_ZTISt8bad_cast} PR119734 { target offload_target_nvptx xfail
*-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_nvptx
} } */
#include "../libgomp.oacc-c++/exceptions-bad_cast-2.C"
-/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
- { dg-output {caught 'std::bad_cast'[\r\n]+} } */
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ { dg-output {.*caught 'std::bad_cast'[\r\n]+} { target { ! offload_device }
} }
+ For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+ TODO For GCN, nvptx offload execution, this currently doesn't 'abort' due to
+ the 'std::bad_cast' exception, but rather due to SIGSEGV in 'dynamic_cast';
+ PR119692.
+
+ For GCN, nvptx offload execution, there is no 'catch'ing; any exception is
fatal.
+ { dg-shouldfail {'MyException' exception} { offload_device } } */
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C
index a588016210c4..3cdedf48529e 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C
@@ -1,10 +1,11 @@
-/* Exception handling constructs in dead code. */
+/* Exception handling constructs in dead code,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_amdgcn } }
{ dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -O0 } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -17,7 +18,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized
} }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1
optimized } }
- Given '-O0', offload compilation fails:
+ Given '-O0' and '-foffload-options=-mno-fake-exceptions', offload
compilation fails:
{ dg-regexp {[^\r\n]+: In function 'f':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C
index 0358c687d205..ef996cfff7fa 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C
@@ -1,10 +1,11 @@
-/* Exception handling constructs in dead code. */
+/* Exception handling constructs in dead code,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_nvptx } }
{ dg-additional-options -foffload=nvptx-none } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -O0 } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -17,7 +18,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized
} }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1
optimized } }
- Given '-O0', offload compilation fails:
+ Given '-O0' and '-foffload-options=-mno-fake-exceptions', offload
compilation fails:
{ dg-regexp {[^\r\n]+: In function 'f':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C
index 8e5ca5cc0b24..a73e7f897be7 100644
--- a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C
@@ -2,16 +2,17 @@
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
-/* { dg-additional-options -foffload=disable }
- Offloading compilation not yet supported; see
- 'target-exceptions-pr118794-1-offload-sorry-GCN.C',
- 'target-exceptions-pr118794-1-offload-sorry-nvptx.C'. */
/* { dg-additional-options -O0 } */
-/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C',
'../../../gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C'. */
+/* Help nvptx offloading overcome a code generation issue;
+ PR106445, PR118518. */
+#define ALWAYS_INLINE __attribute__((always_inline))
+
#pragma omp begin declare target
bool ok = false;
@@ -19,10 +20,12 @@ bool ok = false;
template <typename T>
struct C
{
+ ALWAYS_INLINE
C()
{
ok = true;
}
+ ALWAYS_INLINE
C(int) {};
~C() {};
@@ -55,4 +58,6 @@ int main()
/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes
'__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f':
{ dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1
optimized } }
- { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } } */
+ { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, }
1 optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call
<__builtin_eh_pointer, } 1 optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call
<__builtin_unwind_resume, } 1 optimized } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C
new file mode 100644
index 000000000000..b7a311d36f54
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C
@@ -0,0 +1,25 @@
+/* 'throw' in OpenMP 'target' region, caught. */
+
+/* { dg-additional-options -O0 } */
+/* { dg-require-effective-target exceptions }
+ { dg-additional-options -fexceptions } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+/* { dg-bogus {undefined symbol: typeinfo name for MyException} PR119806 {
target offload_target_amdgcn xfail *-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_amdgcn
} } */
+/* { dg-bogus {Initial value type mismatch} PR119806 { target
offload_target_nvptx xfail *-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_nvptx
} } */
+
+#include "target-exceptions-throw-2.C"
+
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
+ { dg-final { scan-offload-tree-dump-times {gimple_call
<__cxa_allocate_exception, } 1 optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1
optimized } }
+ { dg-output {.*caught 'MyException'[\r\n]+} { target { ! offload_device } }
}
+ For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+ For GCN, nvptx offload execution, there is no 'catch'ing; any exception is
fatal.
+ { dg-shouldfail {'MyException' exception} { offload_device } } */
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
index 484e3f8bd33b..9905b1f21abf 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
@@ -1,10 +1,11 @@
-/* 'throw' in OpenMP 'target' region, caught. */
+/* 'throw' in OpenMP 'target' region, caught,
-foffload-options=-mno-fake-exceptions. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_amdgcn } }
{ dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -14,7 +15,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1
optimized } }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git
a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
index 8ef16020453c..da267d620724 100644
---
a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
+++
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
@@ -1,10 +1,11 @@
-/* 'throw' in OpenMP 'target' region, caught. */
+/* 'throw' in OpenMP 'target' region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target offload_target_nvptx } }
{ dg-additional-options -foffload=nvptx-none } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -14,7 +15,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1
optimized } }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
index f66378ae8653..e85e6c36c422 100644
--- a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
@@ -2,14 +2,22 @@
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
-/* { dg-additional-options -foffload=disable }
- Offloading compilation not yet supported; see
- 'target-exceptions-throw-2-offload-sorry-GCN.C',
- 'target-exceptions-throw-2-offload-sorry-nvptx.C'. */
-/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+/* { dg-bogus {Size expression must be absolute\.} PR119737 { target
offload_target_amdgcn xfail *-*-* } 0 }
+ { dg-ice PR119737 { offload_target_amdgcn } }
+ { dg-excess-errors {'mkoffload' failures etc.} { xfail
offload_target_amdgcn } } */
#include "../libgomp.oacc-c++/exceptions-throw-2.C"
-/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
- { dg-output {caught 'MyException'[\r\n]+} } */
+ { dg-final { scan-offload-tree-dump-times {gimple_call
<__cxa_allocate_exception, } 1 optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1
optimized } }
+ { dg-output {.*caught 'MyException'[\r\n]+} { target { ! offload_device } }
}
+ For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+ For GCN, nvptx offload execution, there is no 'catch'ing; any exception is
fatal.
+ { dg-shouldfail {'MyException' exception} { offload_device } } */
diff --git
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C
index ba6812f9b078..82609663d498 100644
---
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C
+++
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C
@@ -1,9 +1,10 @@
-/* 'std::bad_cast' exception in OpenACC compute region, caught. */
+/* 'std::bad_cast' exception in OpenACC compute region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target openacc_radeon_accel_selected } } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -11,7 +12,7 @@
/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C
index 563f31727ad3..86d3f6c0a08f 100644
---
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C
+++
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C
@@ -1,9 +1,10 @@
-/* 'std::bad_cast' exception in OpenACC compute region, caught. */
+/* 'std::bad_cast' exception in OpenACC compute region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target openacc_nvidia_accel_selected } } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -13,7 +14,7 @@
/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C
index 5f79d63f3826..24399eff978b 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C
@@ -2,17 +2,17 @@
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
-/* { dg-skip-if {} { ! openacc_host_selected } }
- Offloading compilation not yet supported; see
- 'exceptions-bad_cast-2-offload-sorry-GCN.C',
- 'exceptions-bad_cast-2-offload-sorry-nvptx.C'. */
-/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+/* { dg-bogus {_ZTISt8bad_cast} PR119734 { target
openacc_nvidia_accel_selected xfail *-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail
openacc_nvidia_accel_selected } } */
/* See also '../libgomp.c++/target-exceptions-bad_cast-2.C'. */
/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C',
'../../../gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C'. */
+#include <iostream>
#include <typeinfo>
struct C1
@@ -27,6 +27,7 @@ struct C2 : C1
int main()
{
+ std::cerr << "CheCKpOInT\n";
#pragma omp target
#pragma acc serial
{
@@ -44,5 +45,16 @@ int main()
}
}
-/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
- { dg-output {caught 'std::bad_cast'[\r\n]+} } */
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1
optimized } }
+ { dg-output {.*caught 'std::bad_cast'[\r\n]+} { target
openacc_host_selected } }
+ For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+ TODO For GCN, nvptx offload execution, this currently doesn't 'abort' due to
+ the 'std::bad_cast' exception, but rather due to SIGSEGV in 'dynamic_cast';
+ PR119692.
+
+ For GCN, nvptx offload execution, there is no 'catch'ing; any exception is
fatal.
+ { dg-shouldfail {'std::bad_cast' exception} { ! openacc_host_selected } } */
diff --git
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
index 01c6da97ff14..40be8371671d 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
@@ -1,9 +1,10 @@
-/* 'throw' in OpenACC compute region, caught. */
+/* 'throw' in OpenACC compute region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target openacc_radeon_accel_selected } } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -13,7 +14,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1
optimized } }
{ dg-final { only_for_offload_target amdgcn-amdhsa
scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
index e66b6c59f9d3..94614551a566 100644
---
a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
+++
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
@@ -1,9 +1,10 @@
-/* 'throw' in OpenACC compute region, caught. */
+/* 'throw' in OpenACC compute region, caught,
'-foffload-options=-mno-fake-exceptions'. */
/* As this test case involves an expected offload compilation failure, we have
to handle each offload target individually.
{ dg-do link { target openacc_nvidia_accel_selected } } */
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
+/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */
/* { dg-additional-options -fdump-tree-optimized-raw }
{ dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
@@ -15,7 +16,7 @@
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1
optimized } }
{ dg-final { only_for_offload_target nvptx-none
scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } }
- Offload compilation fails:
+ Given '-foffload-options=-mno-fake-exceptions', offload compilation fails:
{ dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry,
unimplemented: exception handling not supported[\r\n]+)+} }
(Note, using 'dg-regexp' instead of 'dg-message', as the former runs before
the auto-mark-UNSUPPORTED.)
{ dg-excess-errors {'mkoffload' failure etc.} } */
diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C
b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C
index 78ce0a3f946b..f6dc9702b690 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C
@@ -2,25 +2,33 @@
/* { dg-require-effective-target exceptions }
{ dg-additional-options -fexceptions } */
-/* { dg-skip-if {} { ! openacc_host_selected } }
- Offloading compilation not yet supported; see
- 'exceptions-throw-2-offload-sorry-GCN.C',
- 'exceptions-throw-2-offload-sorry-nvptx.C'. */
-/* { dg-additional-options -fdump-tree-optimized-raw } */
+/* { dg-additional-options -fdump-tree-optimized-raw }
+ { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+/* { dg-bogus {undefined symbol: typeinfo name for MyException} PR119806 {
target { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } xfail *-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail {
openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } } */
+/* { dg-bogus {Size expression must be absolute\.} PR119737 { target {
openacc_radeon_accel_selected && __OPTIMIZE__ } xfail *-*-* } 0 }
+ { dg-ice PR119737 { openacc_radeon_accel_selected && __OPTIMIZE__ } }
+ { dg-excess-errors {'mkoffload' failures etc.} { xfail {
openacc_radeon_accel_selected && __OPTIMIZE__ } } } */
+/* { dg-bogus {Initial value type mismatch} PR119806 { target {
openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } xfail *-*-* } 0 }
+ { dg-excess-errors {'mkoffload' failure etc.} { xfail {
openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } } */
/* See also '../libgomp.c++/target-exceptions-throw-2.C'. */
/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-throw-2.C',
'../../../gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C'. */
+#include <iostream>
+
class MyException
{
};
int main()
{
+ std::cerr << "CheCKpOInT\n";
#pragma omp target
#pragma acc serial
+ /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target
openacc_nvidia_accel_selected xfail *-*-* } .-1 } */
{
try
{
@@ -34,6 +42,14 @@ int main()
}
}
-/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
+/* { dg-output {CheCKpOInT[\r\n]+} }
+
+ { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, }
1 optimized } }
{ dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized
} }
- { dg-output {caught 'MyException'[\r\n]+} } */
+ { dg-final { scan-offload-tree-dump-times {gimple_call
<__cxa_allocate_exception, } 1 optimized } }
+ { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1
optimized } }
+ { dg-output {.*caught 'MyException'[\r\n]+} { target openacc_host_selected
} }
+ For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+ For GCN, nvptx offload execution, there is no 'catch'ing; any exception is
fatal.
+ { dg-shouldfail {'MyException' exception} { ! openacc_host_selected } } */
