clang/test/SemaCUDA/implicit-device-lambda.cu

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