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.

16 files changed, 140 insertions, 54 deletions

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 ed670a8..93884df 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 8450630..83ec89b 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 26decf6..8861740 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 a588016..3cdedf4 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 0358c68..ef996cf 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 8e5ca5c..a73e7f8 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 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 484e3f8..9905b1f 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 8ef1602..da267d6 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 f66378a..e85e6c3 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 ba6812f..8260966 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 563f317..86d3f6c 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 5f79d63f..24399ef 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 01c6da9..40be837 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 e66b6c5..9461455 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 78ce0a3..f6dc970 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 } } */