109 lines · plain
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 %s3// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only \4// RUN: -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s5 6#include "Inputs/cuda.h"7 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 a22 // 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.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 device64 // function.65}66 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 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#endif90 91 auto f3 = [&] __host__ {};92 f3();93#ifdef __CUDA_ARCH__94 // expected-error@-2 {{reference to __host__ function}}95#endif96 97 auto f4 = [&] __host__ __device__ {};98 f4();99}100 101// The special treatment above only applies to lambdas.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