[PATCH] D42923: [CUDA] Allow external variables in separate compilation
Jonas Hahnfeld via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Mon Feb 5 11:23:39 PST 2018
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>
More information about the cfe-commits
mailing list