diff options
| author | Justin Lebar <jlebar@google.com> | 2016-03-30 23:30:21 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-03-30 23:30:21 +0000 |
| commit | ba122ab42fe54aee3427dc61b765cc8a9dad9d85 (patch) | |
| tree | a8167ca817e1ce33986b9b601c5f725379de603b /clang/lib/Sema/SemaDecl.cpp | |
| parent | 0cda7644301dd0793f796681030b9425b227f157 (diff) | |
| download | bcm5719-llvm-ba122ab42fe54aee3427dc61b765cc8a9dad9d85.tar.gz bcm5719-llvm-ba122ab42fe54aee3427dc61b765cc8a9dad9d85.zip | |
[CUDA] Make unattributed constexpr functions implicitly host+device.
With this patch, by a constexpr function is implicitly host+device
unless:
a) it's a variadic function (variadic functions are not allowed on the
device side), or
b) it's preceeded by a __device__ overload in a system header.
The restriction on overloading __host__ __device__ functions on the
basis of their CUDA attributes remains in place, but we use (b) to allow
us to define __device__ overloads for constexpr functions in cmath,
which would otherwise be __host__ __device__ and thus not overloadable.
You can disable this behavior with -fno-cuda-host-device-constexpr.
Reviewers: tra, rnk, rsmith
Subscribers: cfe-commits
Differential Revision: http://reviews.llvm.org/D18380
llvm-svn: 264964
Diffstat (limited to 'clang/lib/Sema/SemaDecl.cpp')
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 3 |
1 files changed, 3 insertions, 0 deletions
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 66abb655bbc..0d91e976d47 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -8009,6 +8009,9 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC, // Handle attributes. ProcessDeclAttributes(S, NewFD, D); + if (getLangOpts().CUDA) + maybeAddCUDAHostDeviceAttrs(S, NewFD, Previous); + if (getLangOpts().OpenCL) { // OpenCL v1.1 s6.5: Using an address space qualifier in a function return // type declaration will generate a compilation error. |
