From a4005e13f72ece696cce5d190fb73abec116d18a Mon Sep 17 00:00:00 2001 From: Yaxun Liu Date: Sat, 28 Jul 2018 03:05:25 +0000 Subject: [CUDA][HIP] Allow function-scope static const variable CUDA 8.0 E.3.9.4 says: Within the body of a __device__ or __global__ function, only __shared__ variables or variables without any device memory qualifiers may be declared with static storage class. It is unclear how a function-scope non-const static variable without device memory qualifier is implemented, therefore only static const variable without device memory qualifier is allowed, which can be emitted as a global variable in constant address space. Currently clang only allows function-scope static variable with __shared__ qualifier. This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space. Differential Revision: https://reviews.llvm.org/D49931 llvm-svn: 338188 --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++++ 1 file changed, 4 insertions(+) (limited to 'clang/lib/CodeGen/CodeGenModule.cpp') diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ecdf78d..ae2c134 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3176,6 +3176,10 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { return LangAS::cuda_constant; else if (D && D->hasAttr()) return LangAS::cuda_shared; + else if (D && D->hasAttr()) + return LangAS::cuda_device; + else if (D && D->getType().isConstQualified()) + return LangAS::cuda_constant; else return LangAS::cuda_device; } -- cgit v1.1