jlebar created this revision. jlebar added a reviewer: tra. jlebar added a subscriber: cfe-commits.
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. https://reviews.llvm.org/D25166 Files: clang/lib/Sema/SemaDecl.cpp clang/test/CodeGenCUDA/convergent.cu clang/test/CodeGenCUDA/device-var-init.cu clang/test/CodeGenCUDA/nothrow.cu Index: clang/test/CodeGenCUDA/nothrow.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/nothrow.cu @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fexceptions -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \ +// RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s + +// RUN: %clang_cc1 -fcxx-exceptions -fexceptions -triple x86_64-unknown-linux-gnu -emit-llvm \ +// RUN: -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() { 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: } Index: clang/test/CodeGenCUDA/device-var-init.cu =================================================================== --- clang/test/CodeGenCUDA/device-var-init.cu +++ clang/test/CodeGenCUDA/device-var-init.cu @@ -182,9 +182,9 @@ 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) Index: clang/test/CodeGenCUDA/convergent.cu =================================================================== --- clang/test/CodeGenCUDA/convergent.cu +++ clang/test/CodeGenCUDA/convergent.cu @@ -36,8 +36,8 @@ // 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]] = { Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -12074,6 +12074,14 @@ FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation())); } + // CUDA device functions cannot throw. + if (getLangOpts().CUDA && !FD->hasAttr<NoThrowAttr>()) { + CUDAFunctionTarget T = IdentifyCUDATarget(FD); + if (T == CFT_Device || T == CFT_Global || + (getLangOpts().CUDAIsDevice && T == CFT_HostDevice)) + FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation())); + } + IdentifierInfo *Name = FD->getIdentifier(); if (!Name) return;
Index: clang/test/CodeGenCUDA/nothrow.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/nothrow.cu @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fexceptions -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \ +// RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s + +// RUN: %clang_cc1 -fcxx-exceptions -fexceptions -triple x86_64-unknown-linux-gnu -emit-llvm \ +// RUN: -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() { 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: } Index: clang/test/CodeGenCUDA/device-var-init.cu =================================================================== --- clang/test/CodeGenCUDA/device-var-init.cu +++ clang/test/CodeGenCUDA/device-var-init.cu @@ -182,9 +182,9 @@ 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) Index: clang/test/CodeGenCUDA/convergent.cu =================================================================== --- clang/test/CodeGenCUDA/convergent.cu +++ clang/test/CodeGenCUDA/convergent.cu @@ -36,8 +36,8 @@ // 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]] = { Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -12074,6 +12074,14 @@ FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation())); } + // CUDA device functions cannot throw. + if (getLangOpts().CUDA && !FD->hasAttr<NoThrowAttr>()) { + CUDAFunctionTarget T = IdentifyCUDATarget(FD); + if (T == CFT_Device || T == CFT_Global || + (getLangOpts().CUDAIsDevice && T == CFT_HostDevice)) + FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation())); + } + IdentifierInfo *Name = FD->getIdentifier(); if (!Name) return;
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits