diff options
| -rw-r--r-- | clang/include/clang/AST/Decl.h | 5 | ||||
| -rw-r--r-- | clang/lib/AST/Decl.cpp | 17 | ||||
| -rw-r--r-- | clang/lib/Sema/Sema.cpp | 5 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/extern-shared.cu | 27 |
4 files changed, 48 insertions, 6 deletions
diff --git a/clang/include/clang/AST/Decl.h b/clang/include/clang/AST/Decl.h index d16be216807..c10ca9e9670 100644 --- a/clang/include/clang/AST/Decl.h +++ b/clang/include/clang/AST/Decl.h @@ -1456,6 +1456,11 @@ public: void setDescribedVarTemplate(VarTemplateDecl *Template); + // Is this variable known to have a definition somewhere in the complete + // program? This may be true even if the declaration has internal linkage and + // has no definition within this source file. + bool isKnownToBeDefined() const; + // Implement isa/cast/dyncast/etc. static bool classof(const Decl *D) { return classofKind(D->getKind()); } static bool classofKind(Kind K) { return K >= firstVar && K <= lastVar; } diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp index 895d50d1013..08718e0a9e8 100644 --- a/clang/lib/AST/Decl.cpp +++ b/clang/lib/AST/Decl.cpp @@ -2432,6 +2432,23 @@ void VarDecl::setDescribedVarTemplate(VarTemplateDecl *Template) { getASTContext().setTemplateOrSpecializationInfo(this, Template); } +bool VarDecl::isKnownToBeDefined() const { + const auto &LangOpts = getASTContext().getLangOpts(); + // In CUDA mode without relocatable device code, variables of form 'extern + // __shared__ Foo foo[]' are pointers to the base of the GPU core's shared + // memory pool. These are never undefined variables, even if they appear + // inside of an anon namespace or static function. + // + // With CUDA relocatable device code enabled, these variables don't get + // special handling; they're treated like regular extern variables. + if (LangOpts.CUDA && !LangOpts.CUDARelocatableDeviceCode && + hasExternalStorage() && hasAttr<CUDASharedAttr>() && + isa<IncompleteArrayType>(getType())) + return true; + + return hasDefinition(); +} + MemberSpecializationInfo *VarDecl::getMemberSpecializationInfo() const { if (isStaticDataMember()) // FIXME: Remove ? diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index c0997688b2d..aa4f6461b65 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -653,6 +653,11 @@ void Sema::getUndefinedButUsed( !isExternalWithNoLinkageType(VD) && !VD->getMostRecentDecl()->isInline()) continue; + + // Skip VarDecls that lack formal definitions but which we know are in + // fact defined somewhere. + if (VD->isKnownToBeDefined()) + continue; } Undefined.push_back(std::make_pair(ND, UndefinedUse.second)); diff --git a/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu index 2c0c6e0e3df..ec6f3169505 100644 --- a/clang/test/SemaCUDA/extern-shared.cu +++ b/clang/test/SemaCUDA/extern-shared.cu @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s +// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -verify %s +// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s -// RUN: %clang_cc1 -fsyntax-only -fcuda-rdc -verify=rdc %s -// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-rdc -verify=rdc %s -// These declarations are fine in separate compilation mode: -// rdc-no-diagnostics +// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-rdc -verify=rdc %s +// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fcuda-rdc -verify=rdc %s + +// Most of these declarations are fine in separate compilation mode. #include "Inputs/cuda.h" @@ -26,3 +26,18 @@ __host__ __device__ void bar() { extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}} extern __shared__ int global_arr[]; // ok extern __shared__ int global_arr1[1]; // expected-error {{__shared__ variable 'global_arr1' cannot be 'extern'}} + +// Check that, iff we're not in rdc mode, extern __shared__ can appear in an +// anonymous namespace / in a static function without generating a warning +// about a variable with internal linkage but no definition +// (-Wundefined-internal). +namespace { +extern __shared__ int global_arr[]; // rdc-warning {{has internal linkage but is not defined}} +__global__ void in_anon_ns() { + extern __shared__ int local_arr[]; // rdc-warning {{has internal linkage but is not defined}} + + // Touch arrays to generate the warning. + local_arr[0] = 0; // rdc-note {{used here}} + global_arr[0] = 0; // rdc-note {{used here}} +} +} // namespace |
