Canonicalize addrspacecast ConstExpr between different pointer types

[jingyue]

Summary

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.

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

[arsenm]

lib/IR/Constants.cpp

1708–1709

This is a good place to use dyn_cast<VectorType> instead of checking for a vector and then cast<> inside of the getVectorNumElements

[jingyue]

Thanks! I also piggybacked a similar change in InstCombine

[meheff]

Couple nits. Otherwise LGTM.

test/CodeGen/NVPTX/access-non-generic.ll

78

spelling nit: bitcast.

81

Is there also an addrspacecast that you could/should check for? It's more of a question than a request.

[jingyue]

test/CodeGen/NVPTX/access-non-generic.ll

78

Will fix

81

nvptx-favor-non-generic will fold the addrspacecast into the load, so what's after the bitcast will be something like:

load i32 addrspace(3)* %1

I'll check for this addrspace-equipped load in the test as well. Thanks!

[jingyue]

Comments inlined

[eliben]

LGTM

[jingyue]

make the test @ld_int_from_float more explicit

[jingyue]

Fixed an obvious typo

[jingyue]

Committed in r211004

lib/IR/ConstantFold.cpp

Constant *llvm::ConstantFoldCastInstruction(unsigned opc, Constant *V,
 
   // If the cast operand is a constant expression, there's a few things we can
   // do to try to simplify it.
   if (ConstantExpr *CE = dyn_cast<ConstantExpr>(V)) {
     if (CE->isCast()) {
       // 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;
       for (unsigned i = 1, e = CE->getNumOperands(); i != e; ++i)
         if (!CE->getOperand(i)->isNullValue()) {
           isAllNull = false;
           break;
         }

lib/IR/Constants.cpp

Constant *ConstantExpr::getBitCast(Constant *C, Type *DstTy) {
 
   return getFoldedCast(Instruction::BitCast, C, DstTy);
 }
 
 Constant *ConstantExpr::getAddrSpaceCast(Constant *C, Type *DstTy) {
   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);
 }
 
 Constant *ConstantExpr::get(unsigned Opcode, Constant *C1, Constant *C2,
                             unsigned Flags) {
   // Check the operands for consistency first.
   assert(Opcode >= Instruction::BinaryOpsBegin &&
          Opcode <  Instruction::BinaryOpsEnd   &&

lib/Transforms/InstCombine/InstCombineCasts.cpp

Instruction *InstCombiner::visitAddrSpaceCast(AddrSpaceCastInst &CI) {
   // address space. This allows the cast to be exposed to other transforms.
   Value *Src = CI.getOperand(0);
   PointerType *SrcTy = cast<PointerType>(Src->getType()->getScalarType());
   PointerType *DestTy = cast<PointerType>(CI.getType()->getScalarType());
 
   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());
   }
 
   return commonPointerCastTransforms(CI);
 }

test/Assembler/addrspacecast-alias.ll

 ; RUN: llvm-as < %s | llvm-dis | FileCheck %s
 
 ; Test that global aliases are allowed to be constant addrspacecast
 
 @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)*)

test/CodeGen/NVPTX/access-non-generic.ll

; PTX: bar.sync 0;
 
   %sum2 = fadd float %1, %3
   %sum3 = fadd float %sum2, %4
   %sum4 = fadd float %sum3, %6
   %sum5 = fadd float %sum4, %9
   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
 }
 
 declare void @llvm.cuda.syncthreads() #3
 
 attributes #3 = { noduplicate nounwind }
 

test/Other/constant-fold-gep.ll

 @p0 = global [4 x i8] zeroinitializer, align 1
 @p12 = addrspace(12) global [4 x i8] zeroinitializer, align 1
 
 define i8* @different_addrspace() nounwind noinline {
 ; OPT: different_addrspace
   %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 {
 ; OPT: same_addrspace
   %p = getelementptr inbounds i8* bitcast ([4 x i8] * @p0 to i8*), i32 2
   ret i8* %p
 ; OPT: ret i8* getelementptr inbounds ([4 x i8]* @p0, i32 0, i32 2)
 }
 
 ; CHECK: attributes #0 = { nounwind }

test/Transforms/InstCombine/constant-fold-address-space-pointer.ll

 define i32 @constant_through_array_as_ptrs() {
 ; CHECK-LABEL: @constant_through_array_as_ptrs(
 ; CHECK-NEXT: ret i32 34
   %p = load i32 addrspace(1)* addrspace(2)* addrspace(0)* @indirect, align 4
   %a = load i32 addrspace(1)* addrspace(2)* %p, align 4
   %b = load i32 addrspace(1)* %a, align 4
   ret i32 %b
 }
+
+@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
+}