r194765 - Fix test failures after addrspacecast added.

[Matt Arsenault]

Author: arsenm
Date: Thu Nov 14 20:19:52 2013
New Revision: 194765

URL: http://llvm.org/viewvc/llvm-project?rev=194765&view=rev
Log:
Fix test failures after addrspacecast added.

Bitcasts between address spaces are no longer allowed.

Modified:
    cfe/trunk/lib/CodeGen/CodeGenModule.cpp
    cfe/trunk/test/CodeGenCUDA/address-spaces.cu

Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=194765&r1=194764&r2=194765&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Nov 14 20:19:52 2013
@@ -1557,6 +1557,9 @@ CodeGenModule::GetOrCreateLLVMGlobal(Str
       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(Str
   }
 
   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::EmitGlobalVarDefinit
   // 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);
   }

Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=194765&r1=194764&r2=194765&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Thu Nov 14 20:19:52 2013
@@ -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;