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);
}
}
if (AddrSpace != Ty->getAddressSpace())
- return llvm::ConstantExpr::getBitCast(GV, Ty);
- else
- return GV;
+ return llvm::ConstantExpr::getAddrSpaceCast(GV, Ty);
+
+ return GV;
}
// 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);
}
__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;