llvm-project
55 строк · 1.5 Кб
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
6void host() {
7throw NULL;
8try {} catch(void*) {}
9}
10__device__ void device() {
11throw NULL;
12// expected-error@-1 {{cannot use 'throw' in __device__ function}}
13try {} catch(void*) {}
14// expected-error@-1 {{cannot use 'try' in __device__ function}}
15}
16__global__ void kernel() {
17throw NULL;
18// expected-error@-1 {{cannot use 'throw' in __global__ function}}
19try {} 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
26__host__ __device__ void hd1() {
27throw NULL;
28try {} 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.
36inline __host__ __device__ void hd2() {
37throw NULL;
38try {} catch(void*) {}
39}
40void call_hd2() { hd2(); }
41
42// Error, instantiated on device.
43inline __host__ __device__ void hd3() {
44throw NULL;
45try {} 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
52__device__ void call_hd3() { hd3(); }
53#ifdef __CUDA_ARCH__
54// expected-note@-2 {{called by 'call_hd3'}}
55#endif
56