[CUDA] Allow external variables in separate compilation

[Hahnfeld]

Summary

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.

[tra]

LGTM.

test/SemaCUDA/extern-shared.cu

4

Nit. -verify=rdc is somewhat confusing as there's no rdc prefixes in the checks below. Perhaps something along the lines of -verify=there-should-be-no-errors would be more descriptive.

[Hahnfeld]

test/SemaCUDA/extern-shared.cu

4

There is: rdc-no-diagnostics.

But given that you missed it, maybe I should move the comment declarations are fine between RUN lines and no-diagnostics? Don't know if that helps much though...

[tra]

test/SemaCUDA/extern-shared.cu

4

Oh! I did miss it. Never mind then.

lib/Sema/SemaDeclAttr.cpp

 
 static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
                                                  Attr.getName()))
     return;
   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;
   }
   if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
       S.CUDADiagIfHostCode(Attr.getLoc(), diag::err_cuda_host_shared)
           << S.CurrentCUDATarget())
     return;
   D->addAttr(::new (S.Context) CUDASharedAttr(

test/SemaCUDA/extern-shared.cu

 // 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"
 
 __device__ void foo() {
   extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
   extern __shared__ int arr[];  // ok
   extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}}
   extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}}