[OpenMP][SPIRV] Disable exceptions for OpenMP SPIR-V

[sarnex]

More missed target checks.

[github-actions]

🐧 Linux x64 Test Results

• 111356 tests passed
• 4438 tests skipped

[llvmbot]

@llvm/pr-subscribers-clang

Author: Nick Sarnie (sarnex)

Changes

More missed target checks.

---

Full diff: https://github.com/llvm/llvm-project/pull/169094.diff

3 Files Affected:

• (modified) clang/lib/CodeGen/CGException.cpp (+2-2)
• (modified) clang/lib/Frontend/CompilerInvocation.cpp (+1-2)
• (added) clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp (+9)

diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
index f86af4581c345..e9d20672ce185 100644
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -450,7 +450,7 @@ void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
   // 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())) {
+  if (CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU()) {
     EmitTrapCall(llvm::Intrinsic::trap);
     return;
   }
@@ -627,7 +627,7 @@ void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
   // 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()));
+      (CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU());
   if (!IsTargetDevice)
     EnterCXXTryStmt(S);
   EmitStmt(S.getTryBlock());
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 2e027934a8db6..15b44ef1df3da 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -4303,8 +4303,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
 
   // Set the flag to prevent the implementation from emitting device exception
   // handling code for those requiring so.
-  if ((Opts.OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) ||
-      Opts.OpenCLCPlusPlus) {
+  if ((Opts.OpenMPIsTargetDevice && T.isGPU()) || Opts.OpenCLCPlusPlus) {
 
     Opts.Exceptions = 0;
     Opts.CXXExceptions = 0;
diff --git a/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp b/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp
new file mode 100644
index 0000000000000..42f8f3ea70f7d
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fexceptions -fcxx-exceptions -Wno-openmp-target-exception -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -fexceptions -fcxx-exceptions -Wno-openmp-target-exception -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | \
+// RUN: FileCheck -implicit-check-not='{{invoke|throw|cxa}}' %s
+void foo() {
+  // CHECK: call addrspace(9) void @llvm.trap()
+  // CHECK-NEXT: call spir_func addrspace(9) void @__kmpc_target_deinit()
+  #pragma omp target
+  throw "bad";
+}

[llvmbot]

@llvm/pr-subscribers-clang-codegen

Author: Nick Sarnie (sarnex)

Changes

More missed target checks.

---

Full diff: https://github.com/llvm/llvm-project/pull/169094.diff

3 Files Affected:

• (modified) clang/lib/CodeGen/CGException.cpp (+2-2)
• (modified) clang/lib/Frontend/CompilerInvocation.cpp (+1-2)
• (added) clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp (+9)

diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
index f86af4581c345..e9d20672ce185 100644
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -450,7 +450,7 @@ void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
   // 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())) {
+  if (CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU()) {
     EmitTrapCall(llvm::Intrinsic::trap);
     return;
   }
@@ -627,7 +627,7 @@ void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
   // 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()));
+      (CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU());
   if (!IsTargetDevice)
     EnterCXXTryStmt(S);
   EmitStmt(S.getTryBlock());
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 2e027934a8db6..15b44ef1df3da 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -4303,8 +4303,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
 
   // Set the flag to prevent the implementation from emitting device exception
   // handling code for those requiring so.
-  if ((Opts.OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) ||
-      Opts.OpenCLCPlusPlus) {
+  if ((Opts.OpenMPIsTargetDevice && T.isGPU()) || Opts.OpenCLCPlusPlus) {
 
     Opts.Exceptions = 0;
     Opts.CXXExceptions = 0;
diff --git a/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp b/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp
new file mode 100644
index 0000000000000..42f8f3ea70f7d
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_codegen_noexceptions.cpp
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fexceptions -fcxx-exceptions -Wno-openmp-target-exception -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -fexceptions -fcxx-exceptions -Wno-openmp-target-exception -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | \
+// RUN: FileCheck -implicit-check-not='{{invoke|throw|cxa}}' %s
+void foo() {
+  // CHECK: call addrspace(9) void @llvm.trap()
+  // CHECK-NEXT: call spir_func addrspace(9) void @__kmpc_target_deinit()
+  #pragma omp target
+  throw "bad";
+}