Author: jlebar Date: Thu May 17 09:15:07 2018 New Revision: 332621 URL: http://llvm.org/viewvc/llvm-project?rev=332621&view=rev Log: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
Summary: Previously this triggered a -Wundefined-internal warning. But it's not an undefined variable -- any variable of this form is a pointer to the base of GPU core's shared memory. Reviewers: tra Subscribers: sanjoy, rsmith Differential Revision: https://reviews.llvm.org/D46782 Modified: cfe/trunk/include/clang/AST/Decl.h cfe/trunk/lib/AST/Decl.cpp cfe/trunk/lib/Sema/Sema.cpp cfe/trunk/test/SemaCUDA/extern-shared.cu Modified: cfe/trunk/include/clang/AST/Decl.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/Decl.h?rev=332621&r1=332620&r2=332621&view=diff ============================================================================== --- cfe/trunk/include/clang/AST/Decl.h (original) +++ cfe/trunk/include/clang/AST/Decl.h Thu May 17 09:15:07 2018 @@ -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; } Modified: cfe/trunk/lib/AST/Decl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Decl.cpp?rev=332621&r1=332620&r2=332621&view=diff ============================================================================== --- cfe/trunk/lib/AST/Decl.cpp (original) +++ cfe/trunk/lib/AST/Decl.cpp Thu May 17 09:15:07 2018 @@ -2432,6 +2432,23 @@ void VarDecl::setDescribedVarTemplate(Va 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 ? Modified: cfe/trunk/lib/Sema/Sema.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=332621&r1=332620&r2=332621&view=diff ============================================================================== --- cfe/trunk/lib/Sema/Sema.cpp (original) +++ cfe/trunk/lib/Sema/Sema.cpp Thu May 17 09:15:07 2018 @@ -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)); Modified: cfe/trunk/test/SemaCUDA/extern-shared.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/extern-shared.cu?rev=332621&r1=332620&r2=332621&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/extern-shared.cu (original) +++ cfe/trunk/test/SemaCUDA/extern-shared.cu Thu May 17 09:15:07 2018 @@ -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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits