1 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify=dev,expected -fsyntax-only \ 2 // RUN: -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s 3 // RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only \ 4 // RUN: -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s 5 6 #include "Inputs/cuda.h" 7 device_fn()8__device__ void device_fn() { 9 auto f1 = [&] {}; 10 f1(); // implicitly __device__ 11 12 auto f2 = [&] __device__ {}; 13 f2(); 14 15 auto f3 = [&] __host__ {}; 16 f3(); // expected-error {{no matching function}} 17 18 auto f4 = [&] __host__ __device__ {}; 19 f4(); 20 21 // Now do it all again with '()'s in the lambda declarations: This is a 22 // different parse path. 23 auto g1 = [&]() {}; 24 g1(); // implicitly __device__ 25 26 auto g2 = [&]() __device__ {}; 27 g2(); 28 29 auto g3 = [&]() __host__ {}; 30 g3(); // expected-error {{no matching function}} 31 32 auto g4 = [&]() __host__ __device__ {}; 33 g4(); 34 35 // Once more, with the '()'s in a different place. 36 auto h1 = [&]() {}; 37 h1(); // implicitly __device__ 38 39 auto h2 = [&] __device__ () {}; 40 h2(); 41 42 auto h3 = [&] __host__ () {}; 43 h3(); // expected-error {{no matching function}} 44 45 auto h4 = [&] __host__ __device__ () {}; 46 h4(); 47 } 48 49 // Behaves identically to device_fn. kernel_fn()50__global__ void kernel_fn() { 51 auto f1 = [&] {}; 52 f1(); // implicitly __device__ 53 54 auto f2 = [&] __device__ {}; 55 f2(); 56 57 auto f3 = [&] __host__ {}; 58 f3(); // expected-error {{no matching function}} 59 60 auto f4 = [&] __host__ __device__ {}; 61 f4(); 62 63 // No need to re-test all the parser contortions we test in the device 64 // function. 65 } 66 host_fn()67__host__ void host_fn() { 68 auto f1 = [&] {}; 69 f1(); // implicitly __host__ (i.e., no magic) 70 71 auto f2 = [&] __device__ {}; 72 f2(); // expected-error {{no matching function}} 73 74 auto f3 = [&] __host__ {}; 75 f3(); 76 77 auto f4 = [&] __host__ __device__ {}; 78 f4(); 79 } 80 hd_fn()81__host__ __device__ void hd_fn() { 82 auto f1 = [&] {}; 83 f1(); // implicitly __host__ __device__ 84 85 auto f2 = [&] __device__ {}; 86 f2(); 87 #ifndef __CUDA_ARCH__ 88 // expected-error@-2 {{reference to __device__ function}} 89 #endif 90 91 auto f3 = [&] __host__ {}; 92 f3(); 93 #ifdef __CUDA_ARCH__ 94 // expected-error@-2 {{reference to __host__ function}} 95 #endif 96 97 auto f4 = [&] __host__ __device__ {}; 98 f4(); 99 } 100 101 // The special treatment above only applies to lambdas. foo()102__device__ void foo() { 103 struct X { 104 void foo() {} 105 }; 106 X x; 107 x.foo(); // dev-error {{reference to __host__ function 'foo' in __device__ function}} 108 } 109