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

diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index da54de0bed8868cddc657ccb7f9ca44edc0eb413..f5fdfdb64948ff425a21afa9f1ad80b32b477eb5 100644 (file)
--- 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<llvm::ConstantExpr>(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 9df7e3f4d262ac04d77ad349204987c8893e034c..04344526f40e17443c184aac536b8dd4b8f28537 100644 (file)
--- 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;