From 1daf57049878a79b8bb40fb57ee2370f3cd4d4e7 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tschwinge@baylibre.com>
Date: Thu, 27 Mar 2025 14:46:20 +0100
Subject: [PATCH] Add 'throw', caught test cases for GCN, nvptx target and
 OpenACC, OpenMP 'target' offloading

	gcc/testsuite/
	* g++.target/gcn/exceptions-throw-2.C: New.
	* g++.target/nvptx/exceptions-throw-2.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-throw-2.C: New.
	* 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-throw-2.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.
---
 .../g++.target/gcn/exceptions-throw-2.C       | 13 +++++++
 .../g++.target/nvptx/exceptions-throw-2.C     | 13 +++++++
 ...get-exceptions-throw-2-offload-sorry-GCN.C | 20 ++++++++++
 ...t-exceptions-throw-2-offload-sorry-nvptx.C | 20 ++++++++++
 .../libgomp.c++/target-exceptions-throw-2.C   | 15 +++++++
 .../exceptions-throw-2-offload-sorry-GCN.C    | 19 +++++++++
 .../exceptions-throw-2-offload-sorry-nvptx.C  | 21 ++++++++++
 .../libgomp.oacc-c++/exceptions-throw-2.C     | 39 +++++++++++++++++++
 8 files changed, 160 insertions(+)
 create mode 100644 gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
 create mode 100644 gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
 create mode 100644 libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
 create mode 100644 libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C

diff --git a/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C b/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
new file mode 100644
index 00000000000..0ae64eebd47
--- /dev/null
+++ b/gcc/testsuite/g++.target/gcn/exceptions-throw-2.C
@@ -0,0 +1,13 @@
+/* 'throw', caught.  */
+
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++ standard is sufficient.  */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C"
+
+/* { 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 } }
+   Compilation fails:
+   { dg-regexp {[^\r\n]+: In function 'int 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.)  */
diff --git a/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
new file mode 100644
index 00000000000..0ae64eebd47
--- /dev/null
+++ b/gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C
@@ -0,0 +1,13 @@
+/* 'throw', caught.  */
+
+/* Via the magic string "-std=*++" indicate that testing one (the default) C++ standard is sufficient.  */
+/* { dg-additional-options -fexceptions } */
+/* { dg-additional-options -fdump-tree-optimized-raw } */
+
+#include "../../../../libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C"
+
+/* { 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 } }
+   Compilation fails:
+   { dg-regexp {[^\r\n]+: In function 'int 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.)  */
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
new file mode 100644
index 00000000000..484e3f8bd33
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C
@@ -0,0 +1,20 @@
+/* 'throw' in OpenMP 'target' region, caught.  */
+
+/* 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 -fdump-tree-optimized-raw }
+   { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+
+#include "target-exceptions-throw-2.C"
+
+/* { 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 { 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:
+   { 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
new file mode 100644
index 00000000000..8ef16020453
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C
@@ -0,0 +1,20 @@
+/* 'throw' in OpenMP 'target' region, caught.  */
+
+/* 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 -fdump-tree-optimized-raw }
+   { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+
+#include "target-exceptions-throw-2.C"
+
+/* { 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 { 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:
+   { 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
new file mode 100644
index 00000000000..f66378ae865
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C
@@ -0,0 +1,15 @@
+/* 'throw' in OpenMP 'target' region, caught.  */
+
+/* { 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 } */
+
+#include "../libgomp.oacc-c++/exceptions-throw-2.C"
+
+/* { 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]+} } */
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
new file mode 100644
index 00000000000..01c6da97ff1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C
@@ -0,0 +1,19 @@
+/* 'throw' in OpenACC compute region, caught.  */
+
+/* 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 -fdump-tree-optimized-raw }
+   { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+
+#include "exceptions-throw-2.C"
+
+/* { 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 { 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:
+   { 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
new file mode 100644
index 00000000000..e66b6c59f9d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C
@@ -0,0 +1,21 @@
+/* 'throw' in OpenACC compute region, caught.  */
+
+/* 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 -fdump-tree-optimized-raw }
+   { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */
+
+#include "exceptions-throw-2.C"
+
+/* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } 0 } */
+
+/* { 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 { 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:
+   { 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
new file mode 100644
index 00000000000..78ce0a3f946
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C
@@ -0,0 +1,39 @@
+/* 'throw' in OpenACC compute region, caught.  */
+
+/* { 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 } */
+
+/* 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'.  */
+
+class MyException
+{
+};
+
+int main()
+{
+#pragma omp target
+#pragma acc serial
+  {
+    try
+      {
+	MyException e1;
+	throw e1;
+      }
+    catch (const MyException &e)
+      {
+	__builtin_printf("caught '%s'\n", "MyException");
+      }
+  }
+}
+
+/* { 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]+} } */