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; }
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; }
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits