56 lines · plain
1// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s2// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s3 4#include "Inputs/cuda.h"5 6void host() {7 throw NULL;8 try {} catch(void*) {}9}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}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 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#endif33}34 35// No error, never instantiated on device.36inline __host__ __device__ void hd2() {37 throw NULL;38 try {} catch(void*) {}39}40void call_hd2() { hd2(); }41 42// Error, instantiated on device.43inline __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#endif50}51 52__device__ void call_hd3() { hd3(); }53#ifdef __CUDA_ARCH__54// expected-note@-2 {{called by 'call_hd3'}}55#endif56