diff options
author | Justin Lebar <jlebar@google.com> | 2016-10-04 23:41:49 +0000 |
---|---|---|
committer | Justin Lebar <jlebar@google.com> | 2016-10-04 23:41:49 +0000 |
commit | e506eb44f35c2080e987d6dafcb5b69af3126227 (patch) | |
tree | 6f93c30a90eba177ed1861f6f33f69e07de44bf4 /test/CodeGenCUDA | |
parent | 8678918d321e6439b535d2c7f4613498b08797f6 (diff) |
[CUDA] Mark device functions as nounwind.
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.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D25166
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283272 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r-- | test/CodeGenCUDA/convergent.cu | 4 | ||||
-rw-r--r-- | test/CodeGenCUDA/device-var-init.cu | 6 | ||||
-rw-r--r-- | test/CodeGenCUDA/nothrow.cu | 39 |
3 files changed, 44 insertions, 5 deletions
diff --git a/test/CodeGenCUDA/convergent.cu b/test/CodeGenCUDA/convergent.cu index 6827c57d29..62818f9e5a 100644 --- a/test/CodeGenCUDA/convergent.cu +++ b/test/CodeGenCUDA/convergent.cu @@ -36,8 +36,8 @@ __host__ __device__ void bar() { // 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]] = { diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu index 6f2d929413..0f4c648134 100644 --- a/test/CodeGenCUDA/device-var-init.cu +++ b/test/CodeGenCUDA/device-var-init.cu @@ -182,9 +182,9 @@ __device__ void df() { 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) diff --git a/test/CodeGenCUDA/nothrow.cu b/test/CodeGenCUDA/nothrow.cu new file mode 100644 index 0000000000..aa5bfcdda6 --- /dev/null +++ b/test/CodeGenCUDA/nothrow.cu @@ -0,0 +1,39 @@ +// 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: } |