blob: f001b57981fa02a4a88e20726d7af626cc6b339e [file] [log] [blame]
Justin Lebar3e6449b2016-10-04 23:41:49 +00001// 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 | \
Justin Lebar4a759ff2016-10-05 00:27:38 +00003// RUN: FileCheck -check-prefix DEVICE %s
Justin Lebar3e6449b2016-10-04 23:41:49 +00004
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]+]]
14void host_fn() { f(); }
15
16// DEVICE: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
17__device__ void foo() {
18 // DEVICE: call void @_Z1fv
19 f();
20}
21
22// DEVICE: define void @_Z12foo_noexceptv() [[DEVICE_ATTR:#[0-9]+]]
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]+]]
30__host__ __device__ void bar() { f(); }
31
32// DEVICE: define void @_Z3bazv() [[DEVICE_ATTR:#[0-9]+]]
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: }