Add 'throw' test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading

	gcc/testsuite/
	* g++.target/gcn/exceptions-throw-1.C: New.
	* g++.target/nvptx/exceptions-throw-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-throw-1.C: New.
	* testsuite/libgomp.c++/target-exceptions-throw-1-O0.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-1.C: Likewise.

3 files changed, 94 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C
new file mode 100644
index 00000000000..00d7c13859d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C
@@ -0,0 +1,23 @@
+/* 'throw' in OpenMP 'target' region.  */
+
+/* { 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 } */
+
+#include "target-exceptions-throw-1.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 } }
+   For host execution, we print something like:
+       terminate called after throwing an instance of 'MyException'
+       Aborted (core dumped)
+   { dg-output {.*MyException} { target { ! offload_device } } }
+   For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+   { dg-shouldfail {'MyException' exception} } */
diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C
new file mode 100644
index 00000000000..246706150cc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C
@@ -0,0 +1,25 @@
+/* 'throw' in OpenMP 'target' region.  */
+
+/* { 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 {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-1.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 } }
+   For host execution, we print something like:
+       terminate called after throwing an instance of 'MyException'
+       Aborted (core dumped)
+   { dg-output {.*MyException} { target { ! offload_device } } }
+   For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+   { dg-shouldfail {'MyException' exception} } */
diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C
new file mode 100644
index 00000000000..f2ef75193c8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C
@@ -0,0 +1,46 @@
+/* 'throw' in OpenACC compute region.  */
+
+/* { 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 {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' failure etc.} { xfail { openacc_radeon_accel_selected && __OPTIMIZE__ } } } */
+
+/* See also '../libgomp.c++/target-exceptions-throw-1.C'.  */
+
+/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-throw-1.C',
+   '../../../gcc/testsuite/g++.target/nvptx/exceptions-throw-1.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 } */
+  {
+    MyException e1;
+    throw e1;
+  }
+}
+
+/* { 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 } }
+   For host execution, we print something like:
+       terminate called after throwing an instance of 'MyException'
+       Aborted (core dumped)
+   { dg-output {.*MyException} { target openacc_host_selected } }
+   For GCN, nvptx offload execution, we don't print anything, but just 'abort'.
+
+   { dg-shouldfail {'MyException' exception} } */