diff options
| -rw-r--r-- | lib/Sema/SemaDecl.cpp | 8 | ||||
| -rw-r--r-- | test/CodeGenCUDA/address-spaces.cu | 8 |
2 files changed, 16 insertions, 0 deletions
diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp index d310abfe4d..2fe4e63dbb 100644 --- a/lib/Sema/SemaDecl.cpp +++ b/lib/Sema/SemaDecl.cpp @@ -4320,6 +4320,14 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, // Handle attributes prior to checking for duplicates in MergeVarDecl ProcessDeclAttributes(S, NewVD, D); + if (getLangOpts().CUDA) { + // CUDA B.2.5: "__shared__ and __constant__ variables have implied static + // storage [duration]." + if (SC == SC_None && S->getFnParent() != 0 && + (NewVD->hasAttr<CUDASharedAttr>() || NewVD->hasAttr<CUDAConstantAttr>())) + NewVD->setStorageClass(SC_Static); + } + // In auto-retain/release, infer strong retension for variables of // retainable type. if (getLangOpts().ObjCAutoRefCount && inferObjCARCLifetime(NewVD)) diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 15e49205b6..9df7e3f4d2 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -24,5 +24,13 @@ __device__ void foo() { static int li; // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li li++; + + __constant__ int lj; + // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj + lj++; + + __shared__ int lk; + // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk + lk++; } |