1 // RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions -fcuda-is-device \
2 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -disable-llvm-passes -o - %s | \
3 // RUN: FileCheck -check-prefix DEVICE %s
4
5 // RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions \
6 // RUN: -triple x86_64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -o - %s | \
7 // RUN: FileCheck -check-prefix HOST %s
8
9 #include "Inputs/cuda.h"
10
11 __host__ __device__ void f();
12
13 // HOST: define void @_Z7host_fnv() [[HOST_ATTR:#[0-9]+]]
host_fn()14 void host_fn() { f(); }
15
16 // DEVICE: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
foo()17 __device__ void foo() {
18 // DEVICE: call void @_Z1fv
19 f();
20 }
21
22 // DEVICE: define void @_Z12foo_noexceptv() [[DEVICE_ATTR:#[0-9]+]]
foo_noexcept()23 __device__ void foo_noexcept() noexcept {
24 // DEVICE: call void @_Z1fv
25 f();
26 }
27
28 // This is nounwind only on the device side.
29 // CHECK: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
bar()30 __host__ __device__ void bar() { f(); }
31
32 // DEVICE: define void @_Z3bazv() [[DEVICE_ATTR:#[0-9]+]]
baz()33 __global__ void baz() { f(); }
34
35 // DEVICE: attributes [[DEVICE_ATTR]] = {
36 // DEVICE-SAME: nounwind
37 // HOST: attributes [[HOST_ATTR]] = {
38 // HOST-NOT: nounwind
39 // HOST-SAME: }
40