From 39766c54d85d4aef2398c1eab284806a08916de0 Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Fri, 30 Sep 2016 17:14:53 +0000 Subject: [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282880 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaCUDA.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) (limited to 'lib/Sema/SemaCUDA.cpp') diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index c75bdc7f59..293baa55e0 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -559,3 +559,22 @@ bool Sema::CheckCUDAVLA(SourceLocation Loc) { } return true; } + +void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) { + if (Method->hasAttr() || Method->hasAttr()) + return; + FunctionDecl *CurFn = dyn_cast(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. +} -- cgit v1.2.1