diff options
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 3 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/extern-shared.cu | 5 |
2 files changed, 7 insertions, 1 deletions
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index cb6b030b712..c36a6d4c6a9 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/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/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu index 9fd7bbba593..2c0c6e0e3df 100644 --- a/clang/test/SemaCUDA/extern-shared.cu +++ b/clang/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() { |
