From: Jonas Hahnfeld <hahnjo@hahnjo.de>
Date: Wed, 14 Feb 2018 16:04:03 +0000 (+0000)
Subject: [CUDA] Allow external variables in separate compilation
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=751ae2ab8b6f9d5f64524ebf9052b6fbbdfdbf37;p=clang

[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 cb6b030b71..c36a6d4c6a 100644
--- 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 9fd7bbba59..2c0c6e0e3d 100644
--- 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() {