[CUDA] Mark device functions as nounwind.

[jlebar]

Summary

This prevents clang from emitting 'invoke's and catch statements.

Things previously mostly worked thanks to TryToMarkNoThrow() in
CodeGenFunction. But this is not a proper IPO, and it doesn't properly
handle cases like mutual recursion.

Fixes bug 30593.

[rnk]

It feels like the right thing is to disable EH in device side compilation, but obviously that won't work because it would reject try/throw in host code. I think instead of doing that, we should make sure that CUDA diagnoses try / catch / throw in device functions, and then do what you've done here.

Also, take a look at CodeGenFunction::getInvokeDestImpl(). I think you should add some checks in there to return nullptr if we're doing a device-side CUDA compilation. That's a much more direct way to ensure we never generate invokes on the device side.

clang/lib/Sema/SemaDecl.cpp

12082 ↗

(On Diff #73222)

Can we use a noexcept exception specification instead of this GCC attribute?

[jlebar]

In D25166#559117, @rnk wrote:

It feels like the right thing is to disable EH in device side compilation, but obviously that won't work because it would reject try/throw in host code.

Exactly.

I think instead of doing that, we should make sure that CUDA diagnoses try / catch / throw in device functions, > and then do what you've done here.

Disallowing try/throw is https://reviews.llvm.org/D25036.

Also, take a look at CodeGenFunction::getInvokeDestImpl(). I think you should add some checks in there to return nullptr if we're doing a device-side CUDA compilation. That's a much more direct way to ensure we never generate invokes on the device side.

Hm...if we did this but didn't mark functions as "this never throws", would that really take care of everything? For example, if we were to call an external function which is itself not marked as noexcept, llvm wouldn't be able to infer that we don't throw.

Are you saying we should do both? I'm happy to do that, but I am not sure, if we keep the nounwind attribute in, whether it's possible to write a testcase.

Can we use a noexcept exception specification instead of this GCC attribute?

I have no attachment to this specific attribute, but wouldn't adding noexcept change the semantics of the program (because traits can detect whether or not noexcept is present)?

[rnk]

In D25166#559168, @jlebar wrote:

Also, take a look at CodeGenFunction::getInvokeDestImpl(). I think you should add some checks in there to return nullptr if we're doing a device-side CUDA compilation. That's a much more direct way to ensure we never generate invokes on the device side.

Hm...if we did this but didn't mark functions as "this never throws", would that really take care of everything? For example, if we were to call an external function which is itself not marked as noexcept, llvm wouldn't be able to infer that we don't throw.

Are you saying we should do both? I'm happy to do that, but I am not sure, if we keep the nounwind attribute in, whether it's possible to write a testcase.

I think you want both: you want all device CUDA function definitions (and declarations!) to be marked nounwind to help the optimizer, but you also want to directly make sure clang codegen knows that exceptions are turned off on the device side by changing getInvokeDestImpl. It's essentially the key place we check LangOpts.Exceptions, so if we can't turn that off, then we probably want to inject our CUDA device checks there.

Can we use a noexcept exception specification instead of this GCC attribute?

I have no attachment to this specific attribute, but wouldn't adding noexcept change the semantics of the program (because traits can detect whether or not noexcept is present)?

Hm, yeah, that's no good because then you'll have host vs. device conflicts in those traits. This should probably be a CodeGen-only change anyway, so you might as well add nounwind to function declarations in CGCall.cpp.

[jlebar]

Move everything into codegen.

[rnk]

lgtm

clang/test/CodeGenCUDA/nothrow.cu

17

I would check for call void @_Z1fv here, and make foo noexcept. The noexcept is useful because it forces us to essentially emit try { ... } catch (...) { std::terminate(); } around the whole function, but we probably allow noexcept.

[jlebar]

Update tests.

clang/lib/CodeGen/CGCall.cpp

void CodeGenModule::ConstructAttributeList(
 
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     // Conservatively, mark all functions and calls in CUDA as convergent
     // (meaning, they may call an intrinsically convergent op, such as
     // __syncthreads(), and so can't have certain optimizations applied around
     // them).  LLVM will remove this attribute where it safely can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
 
+    // Exceptions aren't supported in CUDA device code.
+    FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
+
     // Respect -fcuda-flush-denormals-to-zero.
     if (getLangOpts().CUDADeviceFlushDenormalsToZero)
       FuncAttrs.addAttribute("nvptx-f32ftz", "true");
   }
 
   ClangToLLVMArgMapping IRFunctionArgs(getContext(), FI);
 
   QualType RetTy = FI.getReturnType();

clang/lib/CodeGen/CGException.cpp

llvm::BasicBlock *CodeGenFunction::getInvokeDestImpl() {
   const LangOptions &LO = CGM.getLangOpts();
   if (!LO.Exceptions) {
     if (!LO.Borland && !LO.MicrosoftExt)
       return nullptr;
     if (!currentFunctionUsesSEHTry())
       return nullptr;
   }
 
+  // CUDA device code doesn't have exceptions.
+  if (LO.CUDA && LO.CUDAIsDevice)
+    return nullptr;
+
   // Check the innermost scope for a cached landing pad.  If this is
   // a non-EH cleanup, we'll check enclosing scopes in EmitLandingPad.
   llvm::BasicBlock *LP = EHStack.begin()->getCachedLandingPad();
   if (LP) return LP;
 
   const EHPersonality &Personality = EHPersonality::get(*this);
 
   if (!CurFn->hasPersonalityFn())

clang/test/CodeGenCUDA/convergent.cu

__host__ __device__ void bar() {
   // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]]
   asm volatile ("trap;");
 }
 
 // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // DEVICE: attributes [[BAZ_ATTR]] = {
 // DEVICE-SAME: convergent
 // DEVICE-SAME: }
-// DEVICE: attributes [[CALL_ATTR]] = { convergent }
-// DEVICE: attributes [[ASM_ATTR]] = { convergent
+// DEVICE-DAG: attributes [[CALL_ATTR]] = { convergent
+// DEVICE-DAG: attributes [[ASM_ATTR]] = { convergent
 
 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // HOST: attributes [[BAZ_ATTR]] = {
 // HOST-NOT: convergent
 // NOST-SAME: }

clang/test/CodeGenCUDA/device-var-init.cu

__device__ void df() {
   // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
   static __shared__ ETC s_etc;
   // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 
   // anchor point separating constructors and destructors
   df(); // CHECK: call void @_Z2dfv()
 
   // Verify that we only call non-empty destructors
-  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6
-  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6
-  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6
+  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
+  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
+  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
   // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
   // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
   // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
   // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
   // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
 
   // CHECK-NEXT: ret void
 }
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I

clang/test/CodeGenCUDA/nothrow.cu

+// RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions -fcuda-is-device \
+// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -disable-llvm-passes -o - %s | \
+// RUN  FileCheck -check-prefix DEVICE %s
+
+// RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions \
+// RUN:   -triple x86_64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -o - %s | \
+// RUN:  FileCheck -check-prefix HOST %s
+
+#include "Inputs/cuda.h"
+
+__host__ __device__ void f();
+
+// HOST: define void @_Z7host_fnv() [[HOST_ATTR:#[0-9]+]]
+void host_fn() { f(); }
+
+// DEVICE: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
+__device__ void foo() {
+  // DEVICE: call void @_Z1fv
+  f();
+}
+
+// DEVICE: define void @_Z12foo_noexceptv() [[DEVICE_ATTR:#[0-9]+]]
+__device__ void foo_noexcept() noexcept {
+  // DEVICE: call void @_Z1fv
+  f();
+}
+
+// This is nounwind only on the device side.
+// CHECK: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
+__host__ __device__ void bar() { f(); }
+
+// DEVICE: define void @_Z3bazv() [[DEVICE_ATTR:#[0-9]+]]
+__global__ void baz() { f(); }
+
+// DEVICE: attributes [[DEVICE_ATTR]] = {
+// DEVICE-SAME: nounwind
+// HOST: attributes [[HOST_ATTR]] = {
+// HOST-NOT: nounwind
+// HOST-SAME: }