[llvm] r211004 - Canonicalize addrspacecast ConstExpr between different pointer types
Jingyue Wu
jingyue at google.com
Sun Jun 15 14:40:58 PDT 2014
Author: jingyue
Date: Sun Jun 15 16:40:57 2014
New Revision: 211004
URL: http://llvm.org/viewvc/llvm-project?rev=211004&view=rev
Log:
Canonicalize addrspacecast ConstExpr between different pointer types
As a follow-up to r210375 which canonicalizes addrspacecast
instructions, this patch canonicalizes addrspacecast constant
expressions.
Given clang uses ConstantExpr::getAddrSpaceCast to emit addrspacecast
cosntant expressions, this patch is also a step towards having the
frontend emit canonicalized addrspacecasts.
Piggyback a minor refactor in InstCombineCasts.cpp
Update three affected tests in addrspacecast-alias.ll,
access-non-generic.ll and constant-fold-gep.ll and added one new test in
constant-fold-address-space-pointer.ll
Modified:
llvm/trunk/lib/IR/ConstantFold.cpp
llvm/trunk/lib/IR/Constants.cpp
llvm/trunk/lib/Transforms/InstCombine/InstCombineCasts.cpp
llvm/trunk/test/Assembler/addrspacecast-alias.ll
llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll
llvm/trunk/test/Other/constant-fold-gep.ll
llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll
Modified: llvm/trunk/lib/IR/ConstantFold.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/ConstantFold.cpp?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/lib/IR/ConstantFold.cpp (original)
+++ llvm/trunk/lib/IR/ConstantFold.cpp Sun Jun 15 16:40:57 2014
@@ -529,7 +529,10 @@ Constant *llvm::ConstantFoldCastInstruct
// Try hard to fold cast of cast because they are often eliminable.
if (unsigned newOpc = foldConstantCastPair(opc, CE, DestTy))
return ConstantExpr::getCast(newOpc, CE->getOperand(0), DestTy);
- } else if (CE->getOpcode() == Instruction::GetElementPtr) {
+ } else if (CE->getOpcode() == Instruction::GetElementPtr &&
+ // Do not fold addrspacecast (gep 0, .., 0). It might make the
+ // addrspacecast uncanonicalized.
+ opc != Instruction::AddrSpaceCast) {
// If all of the indexes in the GEP are null values, there is no pointer
// adjustment going on. We might as well cast the source pointer.
bool isAllNull = true;
Modified: llvm/trunk/lib/IR/Constants.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Constants.cpp?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/lib/IR/Constants.cpp (original)
+++ llvm/trunk/lib/IR/Constants.cpp Sun Jun 15 16:40:57 2014
@@ -1698,6 +1698,19 @@ Constant *ConstantExpr::getAddrSpaceCast
assert(CastInst::castIsValid(Instruction::AddrSpaceCast, C, DstTy) &&
"Invalid constantexpr addrspacecast!");
+ // Canonicalize addrspacecasts between different pointer types by first
+ // bitcasting the pointer type and then converting the address space.
+ PointerType *SrcScalarTy = cast<PointerType>(C->getType()->getScalarType());
+ PointerType *DstScalarTy = cast<PointerType>(DstTy->getScalarType());
+ Type *DstElemTy = DstScalarTy->getElementType();
+ if (SrcScalarTy->getElementType() != DstElemTy) {
+ Type *MidTy = PointerType::get(DstElemTy, SrcScalarTy->getAddressSpace());
+ if (VectorType *VT = dyn_cast<VectorType>(DstTy)) {
+ // Handle vectors of pointers.
+ MidTy = VectorType::get(MidTy, VT->getNumElements());
+ }
+ C = getBitCast(C, MidTy);
+ }
return getFoldedCast(Instruction::AddrSpaceCast, C, DstTy);
}
Modified: llvm/trunk/lib/Transforms/InstCombine/InstCombineCasts.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Transforms/InstCombine/InstCombineCasts.cpp?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/lib/Transforms/InstCombine/InstCombineCasts.cpp (original)
+++ llvm/trunk/lib/Transforms/InstCombine/InstCombineCasts.cpp Sun Jun 15 16:40:57 2014
@@ -1919,8 +1919,10 @@ Instruction *InstCombiner::visitAddrSpac
Type *DestElemTy = DestTy->getElementType();
if (SrcTy->getElementType() != DestElemTy) {
Type *MidTy = PointerType::get(DestElemTy, SrcTy->getAddressSpace());
- if (CI.getType()->isVectorTy()) // Handle vectors of pointers.
- MidTy = VectorType::get(MidTy, CI.getType()->getVectorNumElements());
+ if (VectorType *VT = dyn_cast<VectorType>(CI.getType())) {
+ // Handle vectors of pointers.
+ MidTy = VectorType::get(MidTy, VT->getNumElements());
+ }
Value *NewBitCast = Builder->CreateBitCast(Src, MidTy);
return new AddrSpaceCastInst(NewBitCast, CI.getType());
Modified: llvm/trunk/test/Assembler/addrspacecast-alias.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Assembler/addrspacecast-alias.ll?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/test/Assembler/addrspacecast-alias.ll (original)
+++ llvm/trunk/test/Assembler/addrspacecast-alias.ll Sun Jun 15 16:40:57 2014
@@ -4,4 +4,4 @@
@i = internal addrspace(1) global i8 42
@ia = alias internal addrspacecast (i8 addrspace(1)* @i to i8 addrspace(2)* addrspace(3)*)
-; CHECK: @ia = alias internal addrspacecast (i8 addrspace(1)* @i to i8 addrspace(2)* addrspace(3)*)
+; CHECK: @ia = alias internal addrspacecast (i8 addrspace(2)* addrspace(1)* bitcast (i8 addrspace(1)* @i to i8 addrspace(2)* addrspace(1)*) to i8 addrspace(2)* addrspace(3)*)
Modified: llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll Sun Jun 15 16:40:57 2014
@@ -74,13 +74,13 @@ define float @ld_st_shared_f32(i32 %i, f
ret float %sum5
}
-; Verifies nvptx-favor-non-generic keeps addrspacecasts between pointers of
-; different element types.
+; When hoisting an addrspacecast between different pointer types, replace the
+; addrspacecast with a bitcast.
define i32 @ld_int_from_float() {
; IR-LABEL: @ld_int_from_float
-; IR: addrspacecast
+; IR: load i32 addrspace(3)* bitcast (float addrspace(3)* @scalar to i32 addrspace(3)*)
; PTX-LABEL: ld_int_from_float(
-; PTX: cvta.shared.u{{(32|64)}}
+; PTX: ld.shared.u{{(32|64)}}
%1 = load i32* addrspacecast(float addrspace(3)* @scalar to i32*), align 4
ret i32 %1
}
Modified: llvm/trunk/test/Other/constant-fold-gep.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Other/constant-fold-gep.ll?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/test/Other/constant-fold-gep.ll (original)
+++ llvm/trunk/test/Other/constant-fold-gep.ll Sun Jun 15 16:40:57 2014
@@ -457,7 +457,7 @@ define i8* @different_addrspace() nounwi
%p = getelementptr inbounds i8* addrspacecast ([4 x i8] addrspace(12)* @p12 to i8*),
i32 2
ret i8* %p
-; OPT: ret i8* getelementptr (i8* addrspacecast ([4 x i8] addrspace(12)* @p12 to i8*), i32 2)
+; OPT: ret i8* getelementptr (i8* addrspacecast (i8 addrspace(12)* getelementptr inbounds ([4 x i8] addrspace(12)* @p12, i32 0, i32 0) to i8*), i32 2)
}
define i8* @same_addrspace() nounwind noinline {
Modified: llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll?rev=211004&r1=211003&r2=211004&view=diff
==============================================================================
--- llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll (original)
+++ llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll Sun Jun 15 16:40:57 2014
@@ -230,3 +230,13 @@ define i32 @constant_through_array_as_pt
%b = load i32 addrspace(1)* %a, align 4
ret i32 %b
}
+
+ at shared_mem = external addrspace(3) global [0 x i8]
+
+define float @canonicalize_addrspacecast(i32 %i) {
+; CHECK-LABEL: @canonicalize_addrspacecast
+; CHECK-NEXT: getelementptr inbounds float* addrspacecast (float addrspace(3)* bitcast ([0 x i8] addrspace(3)* @shared_mem to float addrspace(3)*) to float*), i32 %i
+ %p = getelementptr inbounds float* addrspacecast ([0 x i8] addrspace(3)* @shared_mem to float*), i32 %i
+ %v = load float* %p
+ ret float %v
+}
More information about the llvm-commits
mailing list