r332621 - [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Thu May 17 09:15:08 PDT 2018
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
More information about the cfe-commits
mailing list