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 --git a/gcc/config/gcn/gcn.md b/gcc/config/gcn/gcn.md
index 695656f692d6eadd6f9aec50468c5f34d2e8794b..e0fb735e5ec9cf0774983de7c9f4239c54205aaf 100644 (file)
--- 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 142b439bbb88aa8600636fc7fbf6b4877ea6bee5..99d6aeb2b30f17cdebe8d1fb58749be065ddef5e 100644 (file)
--- 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 f5b89c9fee7ab66596a7f48cedcaad5daae84af1..b284ff422ee73a5024e237f55b5796b4e1182f37 100644 (file)
--- 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 bdfe7f58b166e268bd259f9c3c84e253bb0856a0..e7ec0ef4f6ac691db9d80d823deba50c9595fcb3 100644 (file)
--- 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 87364bf423c5984bffb1fbb9c75daa6f7e2ff199..28da43ca740b048b1960bc0708507bf8a4c56ae6 100644 (file)
--- 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 3201247c5fb1f5bfc7b53accc06b4481e27fcbba..7c3bd69cc9737e1020feec1759d9b8f8a49f4b24 100644 (file)
--- 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 9be81aecdddfda1b31ad3394cc31cd18fe5d33b5..ce9fbc7312e0bba3fdbad7205964c755962f4110 100644 (file)
--- 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 d5eb9274a62ca8a8864dd199a00e206a89111531..205811c65673928274e90260e6c6c013ac5b21d1 100644 (file)
--- 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 ffe9e19a281e1c0e99a91dbe9898dd356535779b..b047cbed444e8d8c8ec98ae0d150e85f22adda20 100644 (file)
--- 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 (file)
index 0000000..2904188
--- /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 a7013791511d835a9897378d5d92d25486fca76d..20f9d49004cd49101922ba45dd221361182beb2a 100644 (file)
--- 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 (file)
index 0000000..a5f0da2
--- /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 0ae64eebd47861bbf1d5c25729c951050694a24f..671c810a56284f772cfaa451fdcb109f47b52d28 100644 (file)
--- 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 (file)
index 0000000..f1fd505
--- /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 ffe9e19a281e1c0e99a91dbe9898dd356535779b..b047cbed444e8d8c8ec98ae0d150e85f22adda20 100644 (file)
--- 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 (file)
index 0000000..3f40951
--- /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 a7013791511d835a9897378d5d92d25486fca76d..20f9d49004cd49101922ba45dd221361182beb2a 100644 (file)
--- 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 (file)
index 0000000..a5f0da2
--- /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 0ae64eebd47861bbf1d5c25729c951050694a24f..671c810a56284f772cfaa451fdcb109f47b52d28 100644 (file)
--- 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 (file)
index 0000000..f1fd505
--- /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 eae741c55bbbffe6e11e5f2b2957cedb67983652..97e22c035b3d684953d728c4c4eeb928bb9bbd01 100644 (file)
--- 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 eae741c55bbbffe6e11e5f2b2957cedb67983652..97e22c035b3d684953d728c4c4eeb928bb9bbd01 100644 (file)
--- 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 ed670a80fb63ac58d951b7f12d5a3112f79baa31..93884dfdbf84b819285a8ebbb721f0ff9e5729ce 100644 (file)
--- 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 84506304f121e52ba8bc810a1c2b7c8367bdcc3c..83ec89bde8c1deb6fe3ff65f7b87a6ba14461d74 100644 (file)
--- 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 26decf60055d4d898ae835a154386f94730f5963..88617400a723260e9ce1be71ef5a3b4508e8a9a3 100644 (file)
--- 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 a588016210c436845075f728b9187204c662f68b..3cdedf48529e7e000bce27bde6fa1f4e19dc0500 100644 (file)
--- 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 0358c687d20544f5ddbfd946c473ea116ec808e0..ef996cfff7fa5f119e65ce4e008cd0565bd0cc23 100644 (file)
--- 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 8e5ca5cc0b2470050146500bd304942edab7d1b9..a73e7f897be7e42a6a250ca02678f7072e3f2d0d 100644 (file)
--- 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 (file)
index 0000000..b7a311d
--- /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 484e3f8bd33b927fbfbec4aa630c6c13ced02772..9905b1f21abfbbe5fc6a58fc14a8cb01748a0af9 100644 (file)
--- 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 8ef16020453c3d969c9fa54bcec5cb764bc2ca9b..da267d6207240a9000908a2d11cb68f373d59ed2 100644 (file)
--- 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 f66378ae8653dbf07c8e254e42909c8b92f03e8f..e85e6c36c4220f9a4a8de1fe555d1121e19a1362 100644 (file)
--- 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 ba6812f9b0785c03a7bd5911dbbf41332e1dae36..82609663d498ee8a36d78da1ad8349314fd5c123 100644 (file)
--- 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 563f31727ad31997b33b3f6d2509e3b1676a122d..86d3f6c0a08f295f6be2fa74e182c22765a10ee7 100644 (file)
--- 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 5f79d63f3826706caa671b2323c497644322415c..24399eff978b3e8b0d2f275126e7488d90fe1414 100644 (file)
--- 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 01c6da97ff144212a8ddf122eca8dd74933c726c..40be8371671d61e53e2e7e63a0fabcf8d6df2d33 100644 (file)
--- 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 e66b6c59f9d3414529d924c317646bdd9a9e8153..94614551a566747c3d374c8f0df87d15ecc51132 100644 (file)
--- 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 78ce0a3f946b6dc486ad286f290f16fc26c09fcc..f6dc9702b690b28532e3e15f689b026246d0bdbf 100644 (file)
--- 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 } } */