[CUDA] Make lambdas inherit __host__ and __device__ attributes from the scope in which they're created.

Summary: NVCC compat.  Fixes bug 30567.

Reviewers: tra

Subscribers: cfe-commits, rnk

Differential Revision: https://reviews.llvm.org/D25105

llvm-svn: 282880

2 files changed, 113 insertions, 0 deletions

diff --git a/clang/test/SemaCUDA/implicit-device-lambda-hd.cu b/clang/test/SemaCUDA/implicit-device-lambda-hd.cu
new file mode 100644
index 00000000000..6cd0e96af8c
--- /dev/null
+++ b/clang/test/SemaCUDA/implicit-device-lambda-hd.cu
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -verify-ignore-unexpected=note \
+// RUN:   -S -o /dev/null %s
+// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=note \
+// RUN:   -DHOST -S -o /dev/null %s
+#include "Inputs/cuda.h"
+
+__host__ __device__ void hd_fn() {
+  auto f1 = [&] {};
+  f1(); // implicitly __host__ __device__
+
+  auto f2 = [&] __device__ {};
+  f2();
+#ifdef HOST
+  // expected-error@-2 {{reference to __device__ function}}
+#endif
+
+  auto f3 = [&] __host__ {};
+  f3();
+#ifndef HOST
+  // expected-error@-2 {{reference to __host__ function}}
+#endif
+
+  auto f4 = [&] __host__ __device__ {};
+  f4();
+}
+
+
diff --git a/clang/test/SemaCUDA/implicit-device-lambda.cu b/clang/test/SemaCUDA/implicit-device-lambda.cu
new file mode 100644
index 00000000000..d89cdd0b056
--- /dev/null
+++ b/clang/test/SemaCUDA/implicit-device-lambda.cu
@@ -0,0 +1,86 @@
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -fsyntax-only -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+__device__ void device_fn() {
+  auto f1 = [&] {};
+  f1(); // implicitly __device__
+
+  auto f2 = [&] __device__ {};
+  f2();
+
+  auto f3 = [&] __host__ {};
+  f3();  // expected-error {{no matching function}}
+
+  auto f4 = [&] __host__ __device__ {};
+  f4();
+
+  // Now do it all again with '()'s in the lambda declarations: This is a
+  // different parse path.
+  auto g1 = [&]() {};
+  g1(); // implicitly __device__
+
+  auto g2 = [&]() __device__ {};
+  g2();
+
+  auto g3 = [&]() __host__ {};
+  g3();  // expected-error {{no matching function}}
+
+  auto g4 = [&]() __host__ __device__ {};
+  g4();
+
+  // Once more, with the '()'s in a different place.
+  auto h1 = [&]() {};
+  h1(); // implicitly __device__
+
+  auto h2 = [&] __device__ () {};
+  h2();
+
+  auto h3 = [&] __host__ () {};
+  h3();  // expected-error {{no matching function}}
+
+  auto h4 = [&] __host__ __device__ () {};
+  h4();
+}
+
+// Behaves identically to device_fn.
+__global__ void kernel_fn() {
+  auto f1 = [&] {};
+  f1(); // implicitly __device__
+
+  auto f2 = [&] __device__ {};
+  f2();
+
+  auto f3 = [&] __host__ {};
+  f3();  // expected-error {{no matching function}}
+
+  auto f4 = [&] __host__ __device__ {};
+  f4();
+
+  // No need to re-test all the parser contortions we test in the device
+  // function.
+}
+
+__host__ void host_fn() {
+  auto f1 = [&] {};
+  f1(); // implicitly __host__ (i.e., no magic)
+
+  auto f2 = [&] __device__ {};
+  f2();  // expected-error {{no matching function}}
+
+  auto f3 = [&] __host__ {};
+  f3();
+
+  auto f4 = [&] __host__ __device__ {};
+  f4();
+}
+
+// The special treatment above only applies to lambdas.
+__device__ void foo() {
+  struct X {
+    void foo() {}
+  };
+  X x;
+  x.foo(); // expected-error {{reference to __host__ function 'foo' in __device__ function}}
+}