[CUDA] Allow external variables in separate compilation

According to the CUDA Programming Guide this is prohibited in
whole program compilation mode. This makes sense because external
references cannot be satisfied in that mode anyway. However,
such variables are allowed in separate compilation mode which
is a valid use case.

Differential Revision: https://reviews.llvm.org/D42923

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@325136 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index cb6b030b7127597e0960f6666f99d7e3659317f2..c36a6d4c6a9186f2c3e893c1e0c105e54078c82a 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -4112,7 +4112,8 @@ static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   auto *VD = cast<VarDecl>(D);
   // extern __shared__ is only allowed on arrays with no length (e.g.
   // "int x[]").
-  if (VD->hasExternalStorage() && !isa<IncompleteArrayType>(VD->getType())) {
+  if (!S.getLangOpts().CUDARelocatableDeviceCode && VD->hasExternalStorage() &&
+      !isa<IncompleteArrayType>(VD->getType())) {
     S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
     return;
   }

diff --git a/test/SemaCUDA/extern-shared.cu b/test/SemaCUDA/extern-shared.cu
index 9fd7bbba593c959b9712f68edef3eb6ed28e436a..2c0c6e0e3df7873f7c5b9c40f875714cfd7748c8 100644 (file)
--- a/test/SemaCUDA/extern-shared.cu
+++ b/test/SemaCUDA/extern-shared.cu
@@ -1,6 +1,11 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -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
+
 #include "Inputs/cuda.h"
 
 __device__ void foo() {