[cfe-commits] r157167 - in /cfe/trunk: include/clang/Basic/AddressSpaces.h lib/AST/ASTContext.cpp lib/Basic/Targets.cpp lib/CodeGen/CodeGenModule.cpp lib/CodeGen/CodeGenModule.h test/CodeGenCUDA/address-spaces.cu
Peter Collingbourne
peter at pcc.me.uk
Sun May 20 14:08:35 PDT 2012
Author: pcc
Date: Sun May 20 16:08:35 2012
New Revision: 157167
URL: http://llvm.org/viewvc/llvm-project?rev=157167&view=rev
Log:
CUDA: add CodeGen support for global variable address spaces.
Because in CUDA types do not have associated address spaces,
globals are declared in their "native" address space, and accessed
by bitcasting the pointer to address space 0. This relies on address
space 0 being a unified address space.
Added:
cfe/trunk/test/CodeGenCUDA/address-spaces.cu
Modified:
cfe/trunk/include/clang/Basic/AddressSpaces.h
cfe/trunk/lib/AST/ASTContext.cpp
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.h
Modified: cfe/trunk/include/clang/Basic/AddressSpaces.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AddressSpaces.h?rev=157167&r1=157166&r2=157167&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/AddressSpaces.h (original)
+++ cfe/trunk/include/clang/Basic/AddressSpaces.h Sun May 20 16:08:35 2012
@@ -29,6 +29,10 @@
opencl_local,
opencl_constant,
+ cuda_device,
+ cuda_constant,
+ cuda_shared,
+
Last,
Count = Last-Offset
};
Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=157167&r1=157166&r2=157167&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Sun May 20 16:08:35 2012
@@ -206,7 +206,10 @@
static const unsigned FakeAddrSpaceMap[] = {
1, // opencl_global
2, // opencl_local
- 3 // opencl_constant
+ 3, // opencl_constant
+ 4, // cuda_device
+ 5, // cuda_constant
+ 6 // cuda_shared
};
return &FakeAddrSpaceMap;
} else {
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=157167&r1=157166&r2=157167&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Sun May 20 16:08:35 2012
@@ -949,7 +949,10 @@
static const unsigned PTXAddrSpaceMap[] = {
0, // opencl_global
4, // opencl_local
- 1 // opencl_constant
+ 1, // opencl_constant
+ 0, // cuda_device
+ 1, // cuda_constant
+ 4, // cuda_shared
};
class PTXTargetInfo : public TargetInfo {
static const char * const GCCRegNames[];
@@ -3384,7 +3387,10 @@
static const unsigned TCEOpenCLAddrSpaceMap[] = {
3, // opencl_global
4, // opencl_local
- 5 // opencl_constant
+ 5, // opencl_constant
+ 0, // cuda_device
+ 0, // cuda_constant
+ 0 // cuda_shared
};
class TCETargetInfo : public TargetInfo{
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=157167&r1=157166&r2=157167&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Sun May 20 16:08:35 2012
@@ -1172,11 +1172,12 @@
DeferredDecls.erase(DDI);
}
+ unsigned AddrSpace = GetGlobalVarAddressSpace(D, Ty->getAddressSpace());
llvm::GlobalVariable *GV =
new llvm::GlobalVariable(getModule(), Ty->getElementType(), false,
llvm::GlobalValue::ExternalLinkage,
0, MangledName, 0,
- false, Ty->getAddressSpace());
+ false, AddrSpace);
// Handle things which are present even on external declarations.
if (D) {
@@ -1202,7 +1203,10 @@
GV->setThreadLocal(D->isThreadSpecified());
}
- return GV;
+ if (AddrSpace != Ty->getAddressSpace())
+ return llvm::ConstantExpr::getBitCast(GV, Ty);
+ else
+ return GV;
}
@@ -1487,6 +1491,20 @@
return llvmInit;
}
+unsigned CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D,
+ unsigned AddrSpace) {
+ if (LangOpts.CUDA && CodeGenOpts.CUDAIsDevice) {
+ if (D->hasAttr<CUDAConstantAttr>())
+ AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_constant);
+ else if (D->hasAttr<CUDASharedAttr>())
+ AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_shared);
+ else
+ AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_device);
+ }
+
+ return AddrSpace;
+}
+
void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) {
llvm::Constant *Init = 0;
QualType ASTTy = D->getType();
@@ -1566,7 +1584,7 @@
if (GV == 0 ||
GV->getType()->getElementType() != InitType ||
GV->getType()->getAddressSpace() !=
- getContext().getTargetAddressSpace(ASTTy)) {
+ GetGlobalVarAddressSpace(D, getContext().getTargetAddressSpace(ASTTy))) {
// Move the old entry aside so that we'll create a new one.
Entry->setName(StringRef());
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.h?rev=157167&r1=157166&r2=157167&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.h (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.h Sun May 20 16:08:35 2012
@@ -517,6 +517,12 @@
CreateOrReplaceCXXRuntimeVariable(StringRef Name, llvm::Type *Ty,
llvm::GlobalValue::LinkageTypes Linkage);
+ /// GetGlobalVarAddressSpace - Return the address space of the underlying
+ /// global variable for D, as determined by its declaration. Normally this
+ /// is the same as the address space of D's type, but in CUDA, address spaces
+ /// are associated with declarations, not types.
+ unsigned GetGlobalVarAddressSpace(const VarDecl *D, unsigned AddrSpace);
+
/// GetAddrOfGlobalVar - Return the llvm::Constant for the address of the
/// given global variable. If Ty is non-null and if the global doesn't exist,
/// then it will be greated with the specified type instead of whatever the
Added: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=157167&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Sun May 20 16:08:35 2012
@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple ptx32-unknown-unknown | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+// CHECK: @i = global
+__device__ int i;
+
+// CHECK: @j = addrspace(1) global
+__constant__ int j;
+
+// CHECK: @k = addrspace(4) global
+__shared__ int k;
+
+__device__ void foo() {
+ // CHECK: load i32* @i
+ i++;
+
+ // CHECK: load i32* bitcast (i32 addrspace(1)* @j to i32*)
+ j++;
+
+ // CHECK: load i32* bitcast (i32 addrspace(4)* @k to i32*)
+ k++;
+}
+
More information about the cfe-commits
mailing list