[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
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index c75bdc7..293baa5 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -559,3 +559,22 @@
}
return true;
}
+
+void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
+ if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
+ return;
+ FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
+ if (!CurFn)
+ return;
+ CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn);
+ if (Target == CFT_Global || Target == CFT_Device) {
+ Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ } else if (Target == CFT_HostDevice) {
+ Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ }
+
+ // TODO: nvcc doesn't allow you to specify __host__ or __device__ attributes
+ // on lambdas in all contexts -- we should emit a compatibility warning where
+ // we're more permissive.
+}