Index: lib/Sema/SemaDeclAttr.cpp =================================================================== --- lib/Sema/SemaDeclAttr.cpp +++ lib/Sema/SemaDeclAttr.cpp @@ -4112,7 +4112,8 @@ auto *VD = cast(D); // extern __shared__ is only allowed on arrays with no length (e.g. // "int x[]"). - if (VD->hasExternalStorage() && !isa(VD->getType())) { + if (!S.getLangOpts().CUDARelocatableDeviceCode && VD->hasExternalStorage() && + !isa(VD->getType())) { S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD; return; } Index: test/SemaCUDA/extern-shared.cu =================================================================== --- test/SemaCUDA/extern-shared.cu +++ 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() {