[PATCH] D42923: [CUDA] Allow external variables in separate compilation

[Jonas Hahnfeld via Phabricator via cfe-commits]

Hahnfeld created this revision.
Hahnfeld added reviewers: jlebar, tra, hfinkel.
Herald added a subscriber: cfe-commits.

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.


Repository:
  rC Clang

https://reviews.llvm.org/D42923

Files:
  lib/Sema/SemaDeclAttr.cpp
  test/SemaCUDA/extern-shared.cu


Index: test/SemaCUDA/extern-shared.cu
===================================================================
--- test/SemaCUDA/extern-shared.cu
+++ test/SemaCUDA/extern-shared.cu
@@ -1,5 +1,9 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// These declarations are fine in separate compilation mode!
+// RUN: %clang_cc1 -fsyntax-only -fcuda-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-rdc -verify=rdc %s
+// rdc-no-diagnostics
 
 #include "Inputs/cuda.h"
 
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -4112,7 +4112,8 @@
   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;
   }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D42923.132866.patch
Type: text/x-patch
Size: 1148 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20180205/085cb934/attachment.bin>