This revision was automatically updated to reflect the committed changes.
Closed by commit rG0cfc2dba93b1: [OpenMP] Allow exceptions in target regions
when offloading to GPUs (authored by AntonRydahl).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D153924/new/
https://reviews.llvm.org/D153924
Files:
clang/include/clang/Basic/DiagnosticCommonKinds.td
clang/include/clang/Basic/DiagnosticGroups.td
clang/lib/CodeGen/CGException.cpp
clang/lib/Sema/SemaExprCXX.cpp
clang/lib/Sema/SemaStmt.cpp
clang/test/OpenMP/amdgpu_exceptions.cpp
clang/test/OpenMP/amdgpu_throw.cpp
clang/test/OpenMP/amdgpu_throw_trap.cpp
clang/test/OpenMP/amdgpu_try_catch.cpp
clang/test/OpenMP/nvptx_exceptions.cpp
clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
clang/test/OpenMP/nvptx_throw.cpp
clang/test/OpenMP/nvptx_throw_trap.cpp
clang/test/OpenMP/nvptx_try_catch.cpp
clang/test/OpenMP/x86_target_exceptions.cpp
clang/test/OpenMP/x86_target_throw.cpp
clang/test/OpenMP/x86_target_try_catch.cpp
Index: clang/test/OpenMP/x86_target_try_catch.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/x86_target_try_catch.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try {
+ error = 1;
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// expected-no-diagnostics
Index: clang/test/OpenMP/x86_target_throw.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/x86_target_throw.cpp
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+void foo(void) {
+ throw 404;
+}
+#pragma omp end declare target
+// expected-no-diagnostics
Index: clang/test/OpenMP/x86_target_exceptions.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/x86_target_exceptions.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try {
+ throw 404;
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// expected-no-diagnostics
Index: clang/test/OpenMP/nvptx_try_catch.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_try_catch.cpp
@@ -0,0 +1,45 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@36 {{cannot use 'try' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try { // with-warning {{target 'nvptx64' does not support exception handling; 'catch' block is ignored}}
+ error = 1;
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/test/OpenMP/nvptx_throw_trap.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_throw_trap.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=DEVICE %s
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=HOST %s
+// DEVICE: trap;
+// DEVICE-NOT: __cxa_throw
+// HOST: __cxa_throw
+// HOST-NOT: trap;
+#pragma omp declare target
+void foo(void) {
+ throw 404;
+}
+#pragma omp end declare target
Index: clang/test/OpenMP/nvptx_throw.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_throw.cpp
@@ -0,0 +1,38 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+void foo(void) {
+ throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}}
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
===================================================================
--- clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
+++ clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
@@ -34,7 +34,7 @@
#pragma omp declare target
struct S {
int a;
- S(int a) : a(a) { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+ S(int a) : a(a) { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
};
int foo() { return 0; }
@@ -57,7 +57,7 @@
static long aaa = 23;
a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); // expected-note{{called by 'maini1'}}
if (!a)
- throw "Error"; // expected-error {{cannot use 'throw' with exceptions disabled}}
+ throw "Error"; // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
}
} catch(...) {
}
@@ -67,14 +67,14 @@
int baz3() { return 2 + baz2(); }
int baz2() {
#pragma omp target
- try { // expected-error {{cannot use 'try' with exceptions disabled}}
+ try { // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'catch' block is ignored}}
++c;
} catch (...) {
}
return 2 + baz3();
}
-int baz1() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int baz1() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
int foobar1();
int foobar2();
@@ -85,7 +85,7 @@
#pragma omp end declare target
int foobar1() { throw 1; }
-int foobar2() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int foobar2() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
int foobar3();
@@ -95,7 +95,7 @@
int (*D)() = C; // expected-note {{used here}}
// host-note@-1 {{used here}}
#pragma omp end declare target
-int foobar3() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int foobar3() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
// Check no infinite recursion in deferred diagnostic emitter.
long E = (long)&E;
Index: clang/test/OpenMP/nvptx_exceptions.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/nvptx_exceptions.cpp
@@ -0,0 +1,46 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try { // with-warning {{target 'nvptx64' does not support exception handling; 'catch' block is ignored}}
+ throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}}
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/test/OpenMP/amdgpu_try_catch.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/amdgpu_try_catch.cpp
@@ -0,0 +1,45 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@36 {{cannot use 'try' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}}
+ error = 1;
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/test/OpenMP/amdgpu_throw_trap.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/amdgpu_throw_trap.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=DEVICE %s
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=HOST %s
+// DEVICE: s_trap
+// DEVICE-NOT: __cxa_throw
+// HOST: __cxa_throw
+// HOST-NOT: s_trap
+#pragma omp declare target
+void foo(void) {
+ throw 404;
+}
+#pragma omp end declare target
Index: clang/test/OpenMP/amdgpu_throw.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/amdgpu_throw.cpp
@@ -0,0 +1,38 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+void foo(void) {
+ throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}}
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/test/OpenMP/amdgpu_exceptions.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/amdgpu_exceptions.cpp
@@ -0,0 +1,46 @@
+/**
+ * The first four lines test that a warning is produced when enabling
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+ int error = -1;
+ try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}}
+ throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}}
+ }
+ catch (int e){
+ error = e;
+ }
+ return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
Index: clang/lib/Sema/SemaStmt.cpp
===================================================================
--- clang/lib/Sema/SemaStmt.cpp
+++ clang/lib/Sema/SemaStmt.cpp
@@ -4471,13 +4471,22 @@
/// handlers and creates a try statement from them.
StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
ArrayRef<Stmt *> Handlers) {
- // Don't report an error if 'try' is used in system headers.
- if (!getLangOpts().CXXExceptions &&
+ const llvm::Triple &T = Context.getTargetInfo().getTriple();
+ const bool IsOpenMPGPUTarget =
+ getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN());
+ // Don't report an error if 'try' is used in system headers or in an OpenMP
+ // target region compiled for a GPU architecture.
+ if (!IsOpenMPGPUTarget && !getLangOpts().CXXExceptions &&
!getSourceManager().isInSystemHeader(TryLoc) && !getLangOpts().CUDA) {
// Delay error emission for the OpenMP device code.
targetDiag(TryLoc, diag::err_exceptions_disabled) << "try";
}
+ // In OpenMP target regions, we assume that catch is never reached on GPU
+ // targets.
+ if (IsOpenMPGPUTarget)
+ targetDiag(TryLoc, diag::warn_try_not_valid_on_target) << T.str();
+
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
Index: clang/lib/Sema/SemaExprCXX.cpp
===================================================================
--- clang/lib/Sema/SemaExprCXX.cpp
+++ clang/lib/Sema/SemaExprCXX.cpp
@@ -864,13 +864,21 @@
ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex,
bool IsThrownVarInScope) {
- // Don't report an error if 'throw' is used in system headers.
- if (!getLangOpts().CXXExceptions &&
+ const llvm::Triple &T = Context.getTargetInfo().getTriple();
+ const bool IsOpenMPGPUTarget =
+ getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN());
+ // Don't report an error if 'throw' is used in system headers or in an OpenMP
+ // target region compiled for a GPU architecture.
+ if (!IsOpenMPGPUTarget && !getLangOpts().CXXExceptions &&
!getSourceManager().isInSystemHeader(OpLoc) && !getLangOpts().CUDA) {
// Delay error emission for the OpenMP device code.
targetDiag(OpLoc, diag::err_exceptions_disabled) << "throw";
}
+ // In OpenMP target regions, we replace 'throw' with a trap on GPU targets.
+ if (IsOpenMPGPUTarget)
+ targetDiag(OpLoc, diag::warn_throw_not_valid_on_target) << T.str();
+
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
Index: clang/lib/CodeGen/CGException.cpp
===================================================================
--- clang/lib/CodeGen/CGException.cpp
+++ clang/lib/CodeGen/CGException.cpp
@@ -440,6 +440,15 @@
void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
bool KeepInsertionPoint) {
+ // If the exception is being emitted in an OpenMP target region,
+ // and the target is a GPU, we do not support exception handling.
+ // Therefore, we emit a trap which will abort the program, and
+ // prompt a warning indicating that a trap will be emitted.
+ const llvm::Triple &T = Target.getTriple();
+ if (CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) {
+ EmitTrapCall(llvm::Intrinsic::trap);
+ return;
+ }
if (const Expr *SubExpr = E->getSubExpr()) {
QualType ThrowType = SubExpr->getType();
if (ThrowType->isObjCObjectPointerType()) {
@@ -609,9 +618,16 @@
}
void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
- EnterCXXTryStmt(S);
+ const llvm::Triple &T = Target.getTriple();
+ // If we encounter a try statement on in an OpenMP target region offloaded to
+ // a GPU, we treat it as a basic block.
+ const bool IsTargetDevice =
+ (CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN()));
+ if (!IsTargetDevice)
+ EnterCXXTryStmt(S);
EmitStmt(S.getTryBlock());
- ExitCXXTryStmt(S);
+ if (!IsTargetDevice)
+ ExitCXXTryStmt(S);
}
void CodeGenFunction::EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock) {
Index: clang/include/clang/Basic/DiagnosticGroups.td
===================================================================
--- clang/include/clang/Basic/DiagnosticGroups.td
+++ clang/include/clang/Basic/DiagnosticGroups.td
@@ -1292,9 +1292,10 @@
def OpenMPPre51Compat : DiagGroup<"pre-openmp-51-compat">;
def OpenMP51Ext : DiagGroup<"openmp-51-extensions">;
def OpenMPExtensions : DiagGroup<"openmp-extensions">;
+def OpenMPTargetException : DiagGroup<"openmp-target-exception">;
def OpenMP : DiagGroup<"openmp", [
SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget,
- OpenMPMapping, OpenMP51Ext, OpenMPExtensions
+ OpenMPMapping, OpenMP51Ext, OpenMPExtensions, OpenMPTargetException
]>;
// Backend warnings.
Index: clang/include/clang/Basic/DiagnosticCommonKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticCommonKinds.td
+++ clang/include/clang/Basic/DiagnosticCommonKinds.td
@@ -425,4 +425,13 @@
"options %0 and %1 are set to different values">;
def err_opencl_feature_requires : Error<
"feature %0 requires support of %1 feature">;
+
+def warn_throw_not_valid_on_target : Warning<
+ "target '%0' does not support exception handling;"
+ " 'throw' is assumed to be never reached">,
+ InGroup<OpenMPTargetException>;
+def warn_try_not_valid_on_target : Warning<
+ "target '%0' does not support exception handling;"
+ " 'catch' block is ignored">,
+ InGroup<OpenMPTargetException>;
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
