maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA] Allow external variables in separate compilation 

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.

Differential Revision: https://reviews.llvm.org/D42923

llvm-svn: 325136
This commit is contained in:
Jonas Hahnfeld 2018-02-14 16:04:03 +00:00
parent 7e5d525bd5
commit ee47d8cb96
2 changed files with 7 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/Sema/SemaDeclAttr.cpp
View File

@ -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;
}

5
clang/test/SemaCUDA/extern-shared.cu
View File

@ -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() {