1 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \ 2 // RUN: -verify -verify-ignore-unexpected=note 3 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \ 4 // RUN: -verify -verify-ignore-unexpected=note -fopenmp 5 6 // Note: This test won't work with -fsyntax-only, because some of these errors 7 // are emitted during codegen. 8 9 #include "Inputs/cuda.h" 10 device_fn()11__device__ void device_fn() {} 12 // expected-note@-1 5 {{'device_fn' declared here}} 13 14 struct S { SS15 __device__ S() {} 16 // expected-note@-1 2 {{'S' declared here}} ~SS17 __device__ ~S() { device_fn(); } 18 // expected-note@-1 {{'~S' declared here}} 19 int x; 20 }; 21 22 struct T { hdT23 __host__ __device__ void hd() { device_fn(); } 24 // expected-error@-1 {{reference to __device__ function 'device_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() { device_fn(); } 29 30 __host__ __device__ void hd3(); 31 dT32 __device__ void d() {} 33 // expected-note@-1 {{'d' declared here}} 34 }; 35 hd3()36__host__ __device__ void T::hd3() { 37 device_fn(); 38 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} 39 } 40 hd2()41template <typename T> __host__ __device__ void hd2() { device_fn(); } 42 // expected-error@-1 2 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} host_fn()43void host_fn() { hd2<int>(); } 44 hd()45__host__ __device__ void hd() { device_fn(); } 46 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} 47 48 // No error because this is never instantiated. hd3()49template <typename T> __host__ __device__ void hd3() { device_fn(); } 50 local_var()51__host__ __device__ void local_var() { 52 S s; 53 // expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}} 54 } 55 placement_new(char * ptr)56__host__ __device__ void placement_new(char *ptr) { 57 ::new(ptr) S(); 58 // expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}} 59 } 60 explicit_destructor(S * s)61__host__ __device__ void explicit_destructor(S *s) { 62 s->~S(); 63 // expected-error@-1 {{reference to __device__ function '~S' in __host__ __device__ function}} 64 } 65 hd_member_fn()66__host__ __device__ void hd_member_fn() { 67 T t; 68 // Necessary to trigger an error on T::hd. It's (implicitly) inline, so 69 // isn't codegen'ed until we call it. 70 t.hd(); 71 } 72 h_member_fn()73__host__ __device__ void h_member_fn() { 74 T t; 75 t.d(); 76 // expected-error@-1 {{reference to __device__ function 'd' in __host__ __device__ function}} 77 } 78 fn_ptr()79__host__ __device__ void fn_ptr() { 80 auto* ptr = &device_fn; 81 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} 82 } 83 84 template <typename T> fn_ptr_template()85__host__ __device__ void fn_ptr_template() { 86 auto* ptr = &device_fn; // Not an error because the template isn't instantiated. 87 } 88 89 // Launching a kernel from a host function does not result in code generation 90 // for it, so calling HD function which calls a D function should not trigger 91 // errors. hd_func()92static __host__ __device__ void hd_func() { device_fn(); } kernel()93__global__ void kernel() { hd_func(); } host_func(void)94void host_func(void) { kernel<<<1, 1>>>(); } 95 96 // Should allow host function call kernel template with device function argument. 97 __device__ void f(); t()98template<void(*F)()> __global__ void t() { F(); } g()99__host__ void g() { t<f><<<1,1>>>(); } 100