1 // RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -std=c++17 -verify %s 2 3 #include "Inputs/cuda.h" 4 5 // Error, instantiated on device. hasInvalid()6inline __host__ __device__ void hasInvalid() { 7 throw NULL; 8 // expected-error@-1 2{{cannot use 'throw' in __host__ __device__ function}} 9 } 10 hasInvalid2()11inline __host__ __device__ void hasInvalid2() { 12 throw NULL; 13 // expected-error@-1 2{{cannot use 'throw' in __host__ __device__ function}} 14 } 15 hasInvalidDiscarded()16inline __host__ __device__ void hasInvalidDiscarded() { 17 // This is only used in the discarded statements below, so this should not diagnose. 18 throw NULL; 19 } 20 use0()21static __device__ void use0() { 22 hasInvalid(); // expected-note {{called by 'use0'}} 23 hasInvalid(); // expected-note {{called by 'use0'}} 24 25 if constexpr (true) { 26 hasInvalid2(); // expected-note {{called by 'use0'}} 27 } else { 28 hasInvalidDiscarded(); 29 } 30 31 if constexpr (false) { 32 hasInvalidDiscarded(); 33 } else { 34 hasInvalid2(); // expected-note {{called by 'use0'}} 35 } 36 37 if constexpr (false) { 38 hasInvalidDiscarded(); 39 } 40 } 41 42 // To avoid excessive diagnostic messages, deferred diagnostics are only 43 // emitted the first time a function is called. use1()44static __device__ void use1() { 45 use0(); // expected-note 4{{called by 'use1'}} 46 use0(); 47 } 48 use2()49static __device__ void use2() { 50 use1(); // expected-note 4{{called by 'use2'}} 51 use1(); 52 } 53 use3()54static __device__ void use3() { 55 use2(); // expected-note 4{{called by 'use3'}} 56 use2(); 57 } 58 use4()59__global__ void use4() { 60 use3(); // expected-note 4{{called by 'use4'}} 61 use3(); 62 } 63