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

   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
   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.
  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
  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 #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.
 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 }