From 54307dac5910aa02358f217773786d61fd83989d Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Fri, 15 Nov 2013 02:19:52 +0000 Subject: [PATCH] Fix test failures after addrspacecast added. Bitcasts between address spaces are no longer allowed. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@194765 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/CodeGenModule.cpp | 12 ++++++++---- test/CodeGenCUDA/address-spaces.cu | 6 +++--- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp index da54de0bed..f5fdfdb649 100644 --- a/lib/CodeGen/CodeGenModule.cpp +++ b/lib/CodeGen/CodeGenModule.cpp @@ -1557,6 +1557,9 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, return Entry; // Make sure the result is of the correct type. + if (Entry->getType()->getAddressSpace() != Ty->getAddressSpace()) + return llvm::ConstantExpr::getAddrSpaceCast(Entry, Ty); + return llvm::ConstantExpr::getBitCast(Entry, Ty); } @@ -1607,9 +1610,9 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, } if (AddrSpace != Ty->getAddressSpace()) - return llvm::ConstantExpr::getBitCast(GV, Ty); - else - return GV; + return llvm::ConstantExpr::getAddrSpaceCast(GV, Ty); + + return GV; } @@ -1802,7 +1805,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) { // Strip off a bitcast if we got one back. if (llvm::ConstantExpr *CE = dyn_cast(Entry)) { assert(CE->getOpcode() == llvm::Instruction::BitCast || - // all zero index gep. + CE->getOpcode() == llvm::Instruction::AddrSpaceCast || + // All zero index gep. CE->getOpcode() == llvm::Instruction::GetElementPtr); Entry = CE->getOperand(0); } diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 9df7e3f4d2..04344526f4 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -12,13 +12,13 @@ __constant__ int j; __shared__ int k; __device__ void foo() { - // CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*) + // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*) i++; - // CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*) + // CHECK: load i32* addrspacecast (i32 addrspace(4)* @j to i32*) j++; - // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*) + // CHECK: load i32* addrspacecast (i32 addrspace(3)* @k to i32*) k++; static int li; -- 2.40.0