1 // RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s
2 // RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s
3
4 #include "Inputs/cuda.h"
5
host()6 void host() {
7 throw NULL;
8 try {} catch(void*) {}
9 }
device()10 __device__ void device() {
11 throw NULL;
12 // expected-error@-1 {{cannot use 'throw' in __device__ function}}
13 try {} catch(void*) {}
14 // expected-error@-1 {{cannot use 'try' in __device__ function}}
15 }
kernel()16 __global__ void kernel() {
17 throw NULL;
18 // expected-error@-1 {{cannot use 'throw' in __global__ function}}
19 try {} catch(void*) {}
20 // expected-error@-1 {{cannot use 'try' in __global__ function}}
21 }
22
23 // Check that it's an error to use 'try' and 'throw' from a __host__ __device__
24 // function if and only if it's codegen'ed for device.
25
hd1()26 __host__ __device__ void hd1() {
27 throw NULL;
28 try {} catch(void*) {}
29 #ifdef __CUDA_ARCH__
30 // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
31 // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
32 #endif
33 }
34
35 // No error, never instantiated on device.
hd2()36 inline __host__ __device__ void hd2() {
37 throw NULL;
38 try {} catch(void*) {}
39 }
call_hd2()40 void call_hd2() { hd2(); }
41
42 // Error, instantiated on device.
hd3()43 inline __host__ __device__ void hd3() {
44 throw NULL;
45 try {} catch(void*) {}
46 #ifdef __CUDA_ARCH__
47 // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
48 // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
49 #endif
50 }
51
call_hd3()52 __device__ void call_hd3() { hd3(); }
53 #ifdef __CUDA_ARCH__
54 // expected-note@-2 {{called by 'call_hd3'}}
55 #endif
56