1 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ 2 // RUN: -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note 3 4 // Note: This test won't work with -fsyntax-only, because some of these errors 5 // are emitted during codegen. 6 7 #include "Inputs/cuda.h" 8 host_fn()9extern "C" void host_fn() {} 10 // expected-note@-1 7 {{'host_fn' declared here}} 11 12 struct Dummy {}; 13 14 struct S { SS15 S() {} 16 // expected-note@-1 2 {{'S' declared here}} ~SS17 ~S() { host_fn(); } 18 // expected-note@-1 {{'~S' declared here}} 19 int x; 20 }; 21 22 struct T { hdT23 __host__ __device__ void hd() { host_fn(); } 24 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} 25 26 // No error; this is (implicitly) inline and is never called, so isn't 27 // codegen'ed. hd2T28 __host__ __device__ void hd2() { host_fn(); } 29 30 __host__ __device__ void hd3(); 31 hT32 void h() {} 33 // expected-note@-1 {{'h' declared here}} 34 35 void operator+(); 36 // expected-note@-1 {{'operator+' declared here}} 37 operator -T38 void operator-(const T&) {} 39 // expected-note@-1 {{'operator-' declared here}} 40 operator DummyT41 operator Dummy() { return Dummy(); } 42 // expected-note@-1 {{'operator Dummy' declared here}} 43 operator deleteT44 __host__ void operator delete(void *) { host_fn(); }; 45 __device__ void operator delete(void*, __SIZE_TYPE__); 46 }; 47 48 struct U { 49 __device__ void operator delete(void*, __SIZE_TYPE__) = delete; 50 __host__ __device__ void operator delete(void*); 51 }; 52 hd3()53__host__ __device__ void T::hd3() { 54 host_fn(); 55 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} 56 } 57 hd2()58template <typename T> __host__ __device__ void hd2() { host_fn(); } 59 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} kernel()60__global__ void kernel() { hd2<int>(); } 61 hd()62__host__ __device__ void hd() { host_fn(); } 63 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} 64 hd3()65template <typename T> __host__ __device__ void hd3() { host_fn(); } 66 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} device_fn()67__device__ void device_fn() { hd3<int>(); } 68 69 // No error because this is never instantiated. hd4()70template <typename T> __host__ __device__ void hd4() { host_fn(); } 71 local_var()72__host__ __device__ void local_var() { 73 S s; 74 // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}} 75 } 76 placement_new(char * ptr)77__host__ __device__ void placement_new(char *ptr) { 78 ::new(ptr) S(); 79 // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}} 80 } 81 explicit_destructor(S * s)82__host__ __device__ void explicit_destructor(S *s) { 83 s->~S(); 84 // expected-error@-1 {{reference to __host__ function '~S' in __host__ __device__ function}} 85 } 86 class_specific_delete(T * t,U * u)87__host__ __device__ void class_specific_delete(T *t, U *u) { 88 delete t; // ok, call sized device delete even though host has preferable non-sized version 89 delete u; // ok, call non-sized HD delete rather than sized D delete 90 } 91 hd_member_fn()92__host__ __device__ void hd_member_fn() { 93 T t; 94 // Necessary to trigger an error on T::hd. It's (implicitly) inline, so 95 // isn't codegen'ed until we call it. 96 t.hd(); 97 } 98 h_member_fn()99__host__ __device__ void h_member_fn() { 100 T t; 101 t.h(); 102 // expected-error@-1 {{reference to __host__ function 'h' in __host__ __device__ function}} 103 } 104 fn_ptr()105__host__ __device__ void fn_ptr() { 106 auto* ptr = &host_fn; 107 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} 108 } 109 110 template <typename T> fn_ptr_template()111__host__ __device__ void fn_ptr_template() { 112 auto* ptr = &host_fn; // Not an error because the template isn't instantiated. 113 } 114 unaryOp()115__host__ __device__ void unaryOp() { 116 T t; 117 (void) +t; // expected-error {{reference to __host__ function 'operator+' in __host__ __device__ function}} 118 } 119 binaryOp()120__host__ __device__ void binaryOp() { 121 T t; 122 (void) (t - t); // expected-error {{reference to __host__ function 'operator-' in __host__ __device__ function}} 123 } 124 implicitConversion()125__host__ __device__ void implicitConversion() { 126 T t; 127 Dummy d = t; // expected-error {{reference to __host__ function 'operator Dummy' in __host__ __device__ function}} 128 } 129 130 template <typename T> 131 struct TmplStruct { fnTmplStruct132 template <typename U> __host__ __device__ void fn() {} 133 }; 134 135 template <> 136 template <> fn()137__host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); } 138 // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} 139 double_specialization()140__device__ void double_specialization() { TmplStruct<int>().fn<int>(); } 141