diff options
author | Justin Lebar <jlebar@google.com> | 2016-09-30 17:14:53 +0000 |
---|---|---|
committer | Justin Lebar <jlebar@google.com> | 2016-09-30 17:14:53 +0000 |
commit | 39766c54d85d4aef2398c1eab284806a08916de0 (patch) | |
tree | 653a9a03fd806262b957be7749c72d48605e5464 /lib/Sema | |
parent | ebda71892574be571b4855e3698469df30491151 (diff) | |
download | clang-39766c54d85d4aef2398c1eab284806a08916de0.tar.gz |
[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
Diffstat (limited to 'lib/Sema')
-rw-r--r-- | lib/Sema/SemaCUDA.cpp | 19 | ||||
-rw-r--r-- | lib/Sema/SemaLambda.cpp | 7 |
2 files changed, 25 insertions, 1 deletions
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<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. +} diff --git a/lib/Sema/SemaLambda.cpp b/lib/Sema/SemaLambda.cpp index a9462228d6..0de501fc5e 100644 --- a/lib/Sema/SemaLambda.cpp +++ b/lib/Sema/SemaLambda.cpp @@ -886,7 +886,12 @@ void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro, // Attributes on the lambda apply to the method. ProcessDeclAttributes(CurScope, Method, ParamInfo); - + + // CUDA lambdas get implicit attributes based on the scope in which they're + // declared. + if (getLangOpts().CUDA) + CUDASetLambdaAttrs(Method); + // Introduce the function call operator as the current declaration context. PushDeclContext(CurScope, Method); |