diff options
| author | Artem Belevich <tra@google.com> | 2015-08-10 20:57:02 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2015-08-10 20:57:02 +0000 |
| commit | b7e4aab40cd4eb66c07f496c358a0cc0d716ce34 (patch) | |
| tree | 072e0336b260bfe9bbf5d528234057f6a24c4cb7 /clang/lib/Sema/SemaDeclAttr.cpp | |
| parent | a01ff22bb1ef303107c1ce155be5d03ae82f5080 (diff) | |
| download | bcm5719-llvm-b7e4aab40cd4eb66c07f496c358a0cc0d716ce34.tar.gz bcm5719-llvm-b7e4aab40cd4eb66c07f496c358a0cc0d716ce34.zip | |
[CUDA] Add implicit __attribute__((used)) to all __global__ functions.
This allows emitting kernels that were instantiated from the host code
and which would never be explicitly referenced otherwise.
Differential Revision: http://reviews.llvm.org/D11666
llvm-svn: 244501
Diffstat (limited to 'clang/lib/Sema/SemaDeclAttr.cpp')
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 4 |
1 files changed, 4 insertions, 0 deletions
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 31ddd503af9..8e04b69f81a 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3350,6 +3350,10 @@ static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) { D->addAttr(::new (S.Context) CUDAGlobalAttr(Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); + + // Add implicit attribute((used)) so we don't eliminate kernels + // because there is nothing referencing them on device side. + D->addAttr(UsedAttr::CreateImplicit(S.Context)); } static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) { |
