[llvm] r194760 - Add addrspacecast instruction.
Jeroen Ketema
j.ketema at imperial.ac.uk
Mon Nov 18 06:13:18 PST 2013
This commit seems to break compilation of certain CUDA and OpenCL kernels.
For example:
CUDA:
__global__ void foo() {
__shared__ int A[10];
int* p = A;
p[threadIdx.x] = 0;
}
If clang/llvm is complied with debugging enabled I get:
Assertion failed: (CastInst::castIsValid(opc, C, Ty) && "Invalid constantexpr cast!"), function getCast, file llvm/lib/IR/Constants.cpp, line 1448.
OpenCL:
__kernel void foo(int i, __global int *A)
{
__global int *a;
if (i == 0)
a = A;
else
a = NULL;
if (a != NULL)
A[get_global_id(0)] = get_global_id(0);
}
Again I get:
Assertion failed: (CastInst::castIsValid(opc, C, Ty) && "Invalid constantexpr cast!"), function getCast, file llvm/lib/IR/Constants.cpp, line 1448.
Regards,
Jeroen
On Nov 15, 2013, at 1:35 AM, Matthew.Arsenault at amd.com (Matt Arsenault) wrote:
> Author: arsenm
> Date: Thu Nov 14 19:34:59 2013
> New Revision: 194760
>
> URL: http://llvm.org/viewvc/llvm-project?rev=194760&view=rev
> Log:
> Add addrspacecast instruction.
>
> Patch by Michele Scandale!
>
> Modified:
> llvm/trunk/docs/LangRef.rst
> llvm/trunk/include/llvm-c/Core.h
> llvm/trunk/include/llvm/AutoUpgrade.h
> llvm/trunk/include/llvm/Bitcode/LLVMBitCodes.h
> llvm/trunk/include/llvm/CodeGen/ISDOpcodes.h
> llvm/trunk/include/llvm/CodeGen/SelectionDAG.h
> llvm/trunk/include/llvm/CodeGen/SelectionDAGNodes.h
> llvm/trunk/include/llvm/IR/Constants.h
> llvm/trunk/include/llvm/IR/IRBuilder.h
> llvm/trunk/include/llvm/IR/Instruction.def
> llvm/trunk/include/llvm/IR/Instructions.h
> llvm/trunk/include/llvm/InstVisitor.h
> llvm/trunk/include/llvm/Target/TargetLowering.h
> llvm/trunk/lib/Analysis/ConstantFolding.cpp
> llvm/trunk/lib/AsmParser/LLLexer.cpp
> llvm/trunk/lib/AsmParser/LLParser.cpp
> llvm/trunk/lib/AsmParser/LLToken.h
> llvm/trunk/lib/Bitcode/Reader/BitcodeReader.cpp
> llvm/trunk/lib/Bitcode/Writer/BitcodeWriter.cpp
> llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
> llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
> llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h
> llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp
> llvm/trunk/lib/CodeGen/TargetLoweringBase.cpp
> llvm/trunk/lib/IR/AutoUpgrade.cpp
> llvm/trunk/lib/IR/ConstantFold.cpp
> llvm/trunk/lib/IR/Constants.cpp
> llvm/trunk/lib/IR/Core.cpp
> llvm/trunk/lib/IR/Instruction.cpp
> llvm/trunk/lib/IR/Instructions.cpp
> llvm/trunk/lib/IR/Value.cpp
> llvm/trunk/lib/IR/Verifier.cpp
> llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp
> llvm/trunk/lib/Target/X86/X86ISelLowering.cpp
> llvm/trunk/lib/Target/X86/X86ISelLowering.h
> llvm/trunk/test/Assembler/ConstantExprFoldCast.ll
> llvm/trunk/test/Assembler/ConstantExprNoFold.ll
> llvm/trunk/test/CodeGen/X86/atomic-dagsched.ll
> llvm/trunk/test/DebugInfo/X86/dbg-value-isel.ll
> llvm/trunk/test/Feature/newcasts.ll
> llvm/trunk/test/Other/constant-fold-gep.ll
> llvm/trunk/test/Transforms/InstCombine/2009-01-16-PointerAddrSpace.ll
> llvm/trunk/test/Transforms/InstCombine/2012-07-30-addrsp-bitcast.ll
> llvm/trunk/test/Transforms/InstCombine/cast_ptr.ll
> llvm/trunk/test/Transforms/InstCombine/constant-fold-address-space-pointer.ll
> llvm/trunk/test/Transforms/InstCombine/gep-addrspace.ll
> llvm/trunk/test/Transforms/InstCombine/objsize-address-space.ll
> llvm/trunk/unittests/IR/InstructionsTest.cpp
> llvm/trunk/utils/emacs/llvm-mode.el
> llvm/trunk/utils/kate/llvm.xml
> llvm/trunk/utils/vim/llvm.vim
>
> Modified: llvm/trunk/docs/LangRef.rst
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/docs/LangRef.rst?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/docs/LangRef.rst (original)
> +++ llvm/trunk/docs/LangRef.rst Thu Nov 14 19:34:59 2013
> @@ -2382,6 +2382,10 @@ The following is the syntax for constant
> Convert a constant, CST, to another TYPE. The constraints of the
> operands are the same as those for the :ref:`bitcast
> instruction <i_bitcast>`.
> +``addrspacecast (CST to TYPE)``
> + Convert a constant pointer or constant vector of pointer, CST, to another
> + TYPE in a different address space. The constraints of the operands are the
> + same as those for the :ref:`addrspacecast instruction <i_addrspacecast>`.
> ``getelementptr (CSTPTR, IDX0, IDX1, ...)``, ``getelementptr inbounds (CSTPTR, IDX0, IDX1, ...)``
> Perform the :ref:`getelementptr operation <i_getelementptr>` on
> constants. As with the :ref:`getelementptr <i_getelementptr>`
> @@ -5726,9 +5730,9 @@ is always a *no-op cast* because no bits
> conversion. The conversion is done as if the ``value`` had been stored
> to memory and read back as type ``ty2``. Pointer (or vector of
> pointers) types may only be converted to other pointer (or vector of
> -pointers) types with this instruction if the pointer sizes are
> -equal. To convert pointers to other types, use the :ref:`inttoptr
> -<i_inttoptr>` or :ref:`ptrtoint <i_ptrtoint>` instructions first.
> +pointers) types with the same address space through this instruction.
> +To convert pointers to other types, use the :ref:`inttoptr <i_inttoptr>`
> +or :ref:`ptrtoint <i_ptrtoint>` instructions first.
>
> Example:
> """"""""
> @@ -5740,6 +5744,51 @@ Example:
> %Z = bitcast <2 x int> %V to i64; ; yields i64: %V
> %Z = bitcast <2 x i32*> %V to <2 x i64*> ; yields <2 x i64*>
>
> +.. _i_addrspacecast:
> +
> +'``addrspacecast .. to``' Instruction
> +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
> +
> +Syntax:
> +"""""""
> +
> +::
> +
> + <result> = addrspacecast <pty> <ptrval> to <pty2> ; yields pty2
> +
> +Overview:
> +"""""""""
> +
> +The '``addrspacecast``' instruction converts ``ptrval`` from ``pty`` in
> +address space ``n`` to type ``pty2`` in address space ``m``.
> +
> +Arguments:
> +""""""""""
> +
> +The '``addrspacecast``' instruction takes a pointer or vector of pointer value
> +to cast and a pointer type to cast it to, which must have a different
> +address space.
> +
> +Semantics:
> +""""""""""
> +
> +The '``addrspacecast``' instruction converts the pointer value
> +``ptrval`` to type ``pty2``. It can be a *no-op cast* or a complex
> +value modification, depending on the target and the address spaces
> +pair. Pointers conversion within the same address space must be
> +performed with ``bitcast`` instruction. Note that if the address space
> +conversion is legal then both result and operand refer to the same memory
> +location.
> +
> +Example:
> +""""""""
> +
> +.. code-block:: llvm
> +
> + %X = addrspacecast i32* %x to addrspace(1) i32* ; yields addrspace(1) i32*:%x
> + %Y = addrspacecast addrspace(1) <2 x i32>* %y to addrspace(2) i64* ; yields addrspace(2) i32*:%y
> + %Z = addrspacecast <4 x i32*> %z to <4 x float addrspace(3)*> ; yelds <4 x float addrspace(3)*>:%z
> +
> .. _otherops:
>
> Other Operations
>
> Modified: llvm/trunk/include/llvm-c/Core.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm-c/Core.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm-c/Core.h (original)
> +++ llvm/trunk/include/llvm-c/Core.h Thu Nov 14 19:34:59 2013
> @@ -222,6 +222,7 @@ typedef enum {
> LLVMPtrToInt = 39,
> LLVMIntToPtr = 40,
> LLVMBitCast = 41,
> + LLVMAddrSpaceCast = 60,
>
> /* Other Operators */
> LLVMICmp = 42,
> @@ -1159,6 +1160,7 @@ LLVMTypeRef LLVMX86MMXType(void);
> macro(UnaryInstruction) \
> macro(AllocaInst) \
> macro(CastInst) \
> + macro(AddrSpaceCastInst) \
> macro(BitCastInst) \
> macro(FPExtInst) \
> macro(FPToSIInst) \
> @@ -1630,6 +1632,7 @@ LLVMValueRef LLVMConstFPToSI(LLVMValueRe
> LLVMValueRef LLVMConstPtrToInt(LLVMValueRef ConstantVal, LLVMTypeRef ToType);
> LLVMValueRef LLVMConstIntToPtr(LLVMValueRef ConstantVal, LLVMTypeRef ToType);
> LLVMValueRef LLVMConstBitCast(LLVMValueRef ConstantVal, LLVMTypeRef ToType);
> +LLVMValueRef LLVMConstAddrSpaceCast(LLVMValueRef ConstantVal, LLVMTypeRef ToType);
> LLVMValueRef LLVMConstZExtOrBitCast(LLVMValueRef ConstantVal,
> LLVMTypeRef ToType);
> LLVMValueRef LLVMConstSExtOrBitCast(LLVMValueRef ConstantVal,
> @@ -2605,6 +2608,8 @@ LLVMValueRef LLVMBuildIntToPtr(LLVMBuild
> LLVMTypeRef DestTy, const char *Name);
> LLVMValueRef LLVMBuildBitCast(LLVMBuilderRef, LLVMValueRef Val,
> LLVMTypeRef DestTy, const char *Name);
> +LLVMValueRef LLVMBuildAddrSpaceCast(LLVMBuilderRef, LLVMValueRef Val,
> + LLVMTypeRef DestTy, const char *Name);
> LLVMValueRef LLVMBuildZExtOrBitCast(LLVMBuilderRef, LLVMValueRef Val,
> LLVMTypeRef DestTy, const char *Name);
> LLVMValueRef LLVMBuildSExtOrBitCast(LLVMBuilderRef, LLVMValueRef Val,
>
> Modified: llvm/trunk/include/llvm/AutoUpgrade.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/AutoUpgrade.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/AutoUpgrade.h (original)
> +++ llvm/trunk/include/llvm/AutoUpgrade.h Thu Nov 14 19:34:59 2013
> @@ -15,11 +15,14 @@
> #define LLVM_AUTOUPGRADE_H
>
> namespace llvm {
> + class Constant;
> class Module;
> class GlobalVariable;
> class Function;
> class CallInst;
> class Instruction;
> + class Value;
> + class Type;
>
> /// This is a more granular function that simply checks an intrinsic function
> /// for upgrading, and returns true if it requires upgrading. It may return
> @@ -44,6 +47,16 @@ namespace llvm {
> /// If the TBAA tag for the given instruction uses the scalar TBAA format,
> /// we upgrade it to the struct-path aware TBAA format.
> void UpgradeInstWithTBAATag(Instruction *I);
> +
> + /// This is an auto-upgrade for bitcast between pointers with different
> + /// address spaces: the instruction is replaced by a pair ptrtoint+inttoptr.
> + Instruction *UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy,
> + Instruction *&Temp);
> +
> + /// This is an auto-upgrade for bitcast constant expression between pointers
> + /// with different address spaces: the instruction is replaced by a pair
> + /// ptrtoint+inttoptr.
> + Value *UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy);
> } // End llvm namespace
>
> #endif
>
> Modified: llvm/trunk/include/llvm/Bitcode/LLVMBitCodes.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/Bitcode/LLVMBitCodes.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/Bitcode/LLVMBitCodes.h (original)
> +++ llvm/trunk/include/llvm/Bitcode/LLVMBitCodes.h Thu Nov 14 19:34:59 2013
> @@ -194,7 +194,8 @@ namespace bitc {
> CAST_FPEXT = 8,
> CAST_PTRTOINT = 9,
> CAST_INTTOPTR = 10,
> - CAST_BITCAST = 11
> + CAST_BITCAST = 11,
> + CAST_ADDRSPACECAST = 12
> };
>
> /// BinaryOpcodes - These are values used in the bitcode files to encode which
>
> Modified: llvm/trunk/include/llvm/CodeGen/ISDOpcodes.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/CodeGen/ISDOpcodes.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/CodeGen/ISDOpcodes.h (original)
> +++ llvm/trunk/include/llvm/CodeGen/ISDOpcodes.h Thu Nov 14 19:34:59 2013
> @@ -419,6 +419,10 @@ namespace ISD {
> /// getNode().
> BITCAST,
>
> + /// ADDRSPACECAST - This operator converts between pointers of different
> + /// address spaces.
> + ADDRSPACECAST,
> +
> /// CONVERT_RNDSAT - This operator is used to support various conversions
> /// between various types (float, signed, unsigned and vectors of those
> /// types) with rounding and saturation. NOTE: Avoid using this operator as
>
> Modified: llvm/trunk/include/llvm/CodeGen/SelectionDAG.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/CodeGen/SelectionDAG.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/CodeGen/SelectionDAG.h (original)
> +++ llvm/trunk/include/llvm/CodeGen/SelectionDAG.h Thu Nov 14 19:34:59 2013
> @@ -802,6 +802,10 @@ public:
> /// getMDNode - Return an MDNodeSDNode which holds an MDNode.
> SDValue getMDNode(const MDNode *MD);
>
> + /// getAddrSpaceCast - Return an AddrSpaceCastSDNode.
> + SDValue getAddrSpaceCast(SDLoc dl, EVT VT, SDValue Ptr,
> + unsigned SrcAS, unsigned DestAS);
> +
> /// getShiftAmountOperand - Return the specified value casted to
> /// the target's desired shift amount type.
> SDValue getShiftAmountOperand(EVT LHSTy, SDValue Op);
>
> Modified: llvm/trunk/include/llvm/CodeGen/SelectionDAGNodes.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/CodeGen/SelectionDAGNodes.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/CodeGen/SelectionDAGNodes.h (original)
> +++ llvm/trunk/include/llvm/CodeGen/SelectionDAGNodes.h Thu Nov 14 19:34:59 2013
> @@ -950,6 +950,23 @@ public:
> const SDValue &getValue() const { return Op; }
> };
>
> +class AddrSpaceCastSDNode : public UnarySDNode {
> +private:
> + unsigned SrcAddrSpace;
> + unsigned DestAddrSpace;
> +
> +public:
> + AddrSpaceCastSDNode(unsigned Order, DebugLoc dl, EVT VT, SDValue X,
> + unsigned SrcAS, unsigned DestAS);
> +
> + unsigned getSrcAddressSpace() const { return SrcAddrSpace; }
> + unsigned getDestAddressSpace() const { return DestAddrSpace; }
> +
> + static bool classof(const SDNode *N) {
> + return N->getOpcode() == ISD::ADDRSPACECAST;
> + }
> +};
> +
> /// Abstact virtual class for operations for memory operations
> class MemSDNode : public SDNode {
> private:
>
> Modified: llvm/trunk/include/llvm/IR/Constants.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/Constants.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/IR/Constants.h (original)
> +++ llvm/trunk/include/llvm/IR/Constants.h Thu Nov 14 19:34:59 2013
> @@ -862,6 +862,7 @@ public:
> static Constant *getPtrToInt(Constant *C, Type *Ty);
> static Constant *getIntToPtr(Constant *C, Type *Ty);
> static Constant *getBitCast (Constant *C, Type *Ty);
> + static Constant *getAddrSpaceCast (Constant *C, Type *Ty);
>
> static Constant *getNSWNeg(Constant *C) { return getNeg(C, false, true); }
> static Constant *getNUWNeg(Constant *C) { return getNeg(C, true, false); }
>
> Modified: llvm/trunk/include/llvm/IR/IRBuilder.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IRBuilder.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/IR/IRBuilder.h (original)
> +++ llvm/trunk/include/llvm/IR/IRBuilder.h Thu Nov 14 19:34:59 2013
> @@ -1133,6 +1133,10 @@ public:
> const Twine &Name = "") {
> return CreateCast(Instruction::BitCast, V, DestTy, Name);
> }
> + Value *CreateAddrSpaceCast(Value *V, Type *DestTy,
> + const Twine &Name = "") {
> + return CreateCast(Instruction::AddrSpaceCast, V, DestTy, Name);
> + }
> Value *CreateZExtOrBitCast(Value *V, Type *DestTy,
> const Twine &Name = "") {
> if (V->getType() == DestTy)
>
> Modified: llvm/trunk/include/llvm/IR/Instruction.def
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/Instruction.def?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/IR/Instruction.def (original)
> +++ llvm/trunk/include/llvm/IR/Instruction.def Thu Nov 14 19:34:59 2013
> @@ -154,25 +154,26 @@ HANDLE_CAST_INST(41, FPExt , FPExtInst
> HANDLE_CAST_INST(42, PtrToInt, PtrToIntInst) // Pointer -> Integer
> HANDLE_CAST_INST(43, IntToPtr, IntToPtrInst) // Integer -> Pointer
> HANDLE_CAST_INST(44, BitCast , BitCastInst ) // Type cast
> - LAST_CAST_INST(44)
> +HANDLE_CAST_INST(45, AddrSpaceCast, AddrSpaceCastInst) // addrspace cast
> + LAST_CAST_INST(45)
>
> // Other operators...
> - FIRST_OTHER_INST(45)
> -HANDLE_OTHER_INST(45, ICmp , ICmpInst ) // Integer comparison instruction
> -HANDLE_OTHER_INST(46, FCmp , FCmpInst ) // Floating point comparison instr.
> -HANDLE_OTHER_INST(47, PHI , PHINode ) // PHI node instruction
> -HANDLE_OTHER_INST(48, Call , CallInst ) // Call a function
> -HANDLE_OTHER_INST(49, Select , SelectInst ) // select instruction
> -HANDLE_OTHER_INST(50, UserOp1, Instruction) // May be used internally in a pass
> -HANDLE_OTHER_INST(51, UserOp2, Instruction) // Internal to passes only
> -HANDLE_OTHER_INST(52, VAArg , VAArgInst ) // vaarg instruction
> -HANDLE_OTHER_INST(53, ExtractElement, ExtractElementInst)// extract from vector
> -HANDLE_OTHER_INST(54, InsertElement, InsertElementInst) // insert into vector
> -HANDLE_OTHER_INST(55, ShuffleVector, ShuffleVectorInst) // shuffle two vectors.
> -HANDLE_OTHER_INST(56, ExtractValue, ExtractValueInst)// extract from aggregate
> -HANDLE_OTHER_INST(57, InsertValue, InsertValueInst) // insert into aggregate
> -HANDLE_OTHER_INST(58, LandingPad, LandingPadInst) // Landing pad instruction.
> - LAST_OTHER_INST(58)
> + FIRST_OTHER_INST(46)
> +HANDLE_OTHER_INST(46, ICmp , ICmpInst ) // Integer comparison instruction
> +HANDLE_OTHER_INST(47, FCmp , FCmpInst ) // Floating point comparison instr.
> +HANDLE_OTHER_INST(48, PHI , PHINode ) // PHI node instruction
> +HANDLE_OTHER_INST(49, Call , CallInst ) // Call a function
> +HANDLE_OTHER_INST(50, Select , SelectInst ) // select instruction
> +HANDLE_OTHER_INST(51, UserOp1, Instruction) // May be used internally in a pass
> +HANDLE_OTHER_INST(52, UserOp2, Instruction) // Internal to passes only
> +HANDLE_OTHER_INST(53, VAArg , VAArgInst ) // vaarg instruction
> +HANDLE_OTHER_INST(54, ExtractElement, ExtractElementInst)// extract from vector
> +HANDLE_OTHER_INST(55, InsertElement, InsertElementInst) // insert into vector
> +HANDLE_OTHER_INST(56, ShuffleVector, ShuffleVectorInst) // shuffle two vectors.
> +HANDLE_OTHER_INST(57, ExtractValue, ExtractValueInst)// extract from aggregate
> +HANDLE_OTHER_INST(58, InsertValue, InsertValueInst) // insert into aggregate
> +HANDLE_OTHER_INST(59, LandingPad, LandingPadInst) // Landing pad instruction.
> + LAST_OTHER_INST(59)
>
> #undef FIRST_TERM_INST
> #undef HANDLE_TERM_INST
>
> Modified: llvm/trunk/include/llvm/IR/Instructions.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/Instructions.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/IR/Instructions.h (original)
> +++ llvm/trunk/include/llvm/IR/Instructions.h Thu Nov 14 19:34:59 2013
> @@ -3615,6 +3615,43 @@ public:
> }
> };
>
> +//===----------------------------------------------------------------------===//
> +// AddrSpaceCastInst Class
> +//===----------------------------------------------------------------------===//
> +
> +/// \brief This class represents a conversion between pointers from
> +/// one address space to another.
> +class AddrSpaceCastInst : public CastInst {
> +protected:
> + /// \brief Clone an identical AddrSpaceCastInst
> + virtual AddrSpaceCastInst *clone_impl() const;
> +
> +public:
> + /// \brief Constructor with insert-before-instruction semantics
> + AddrSpaceCastInst(
> + Value *S, ///< The value to be casted
> + Type *Ty, ///< The type to casted to
> + const Twine &NameStr = "", ///< A name for the new instruction
> + Instruction *InsertBefore = 0 ///< Where to insert the new instruction
> + );
> +
> + /// \brief Constructor with insert-at-end-of-block semantics
> + AddrSpaceCastInst(
> + Value *S, ///< The value to be casted
> + Type *Ty, ///< The type to casted to
> + const Twine &NameStr, ///< A name for the new instruction
> + BasicBlock *InsertAtEnd ///< The block to insert the instruction into
> + );
> +
> + // Methods for support type inquiry through isa, cast, and dyn_cast:
> + static inline bool classof(const Instruction *I) {
> + return I->getOpcode() == AddrSpaceCast;
> + }
> + static inline bool classof(const Value *V) {
> + return isa<Instruction>(V) && classof(cast<Instruction>(V));
> + }
> +};
> +
> } // End llvm namespace
>
> #endif
>
> Modified: llvm/trunk/include/llvm/InstVisitor.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/InstVisitor.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/InstVisitor.h (original)
> +++ llvm/trunk/include/llvm/InstVisitor.h Thu Nov 14 19:34:59 2013
> @@ -191,6 +191,7 @@ public:
> RetTy visitPtrToIntInst(PtrToIntInst &I) { DELEGATE(CastInst);}
> RetTy visitIntToPtrInst(IntToPtrInst &I) { DELEGATE(CastInst);}
> RetTy visitBitCastInst(BitCastInst &I) { DELEGATE(CastInst);}
> + RetTy visitAddrSpaceCastInst(AddrSpaceCastInst &I) { DELEGATE(CastInst);}
> RetTy visitSelectInst(SelectInst &I) { DELEGATE(Instruction);}
> RetTy visitVAArgInst(VAArgInst &I) { DELEGATE(UnaryInstruction);}
> RetTy visitExtractElementInst(ExtractElementInst &I) { DELEGATE(Instruction);}
>
> Modified: llvm/trunk/include/llvm/Target/TargetLowering.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/Target/TargetLowering.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/include/llvm/Target/TargetLowering.h (original)
> +++ llvm/trunk/include/llvm/Target/TargetLowering.h Thu Nov 14 19:34:59 2013
> @@ -827,6 +827,11 @@ public:
> return 0;
> }
>
> + /// Returns true if a cast between SrcAS and DestAS is a noop.
> + virtual bool isNoopAddrSpaceCast(unsigned SrcAS, unsigned DestAS) const {
> + return false;
> + }
> +
> //===--------------------------------------------------------------------===//
> /// \name Helpers for TargetTransformInfo implementations
> /// @{
>
> Modified: llvm/trunk/lib/Analysis/ConstantFolding.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Analysis/ConstantFolding.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Analysis/ConstantFolding.cpp (original)
> +++ llvm/trunk/lib/Analysis/ConstantFolding.cpp Thu Nov 14 19:34:59 2013
> @@ -663,7 +663,7 @@ static Constant* StripPtrCastKeepAS(Cons
> if (NewPtrTy->getAddressSpace() != OldPtrTy->getAddressSpace()) {
> NewPtrTy = NewPtrTy->getElementType()->getPointerTo(
> OldPtrTy->getAddressSpace());
> - Ptr = ConstantExpr::getBitCast(Ptr, NewPtrTy);
> + Ptr = ConstantExpr::getPointerCast(Ptr, NewPtrTy);
> }
> return Ptr;
> }
> @@ -995,8 +995,9 @@ Constant *llvm::ConstantFoldInstOperands
> return ConstantExpr::getCast(Opcode, Ops[0], DestTy);
> case Instruction::IntToPtr:
> // If the input is a ptrtoint, turn the pair into a ptr to ptr bitcast if
> - // the int size is >= the ptr size. This requires knowing the width of a
> - // pointer, so it can't be done in ConstantExpr::getCast.
> + // the int size is >= the ptr size and the address spaces are the same.
> + // This requires knowing the width of a pointer, so it can't be done in
> + // ConstantExpr::getCast.
> if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Ops[0])) {
> if (TD && CE->getOpcode() == Instruction::PtrToInt) {
> Constant *SrcPtr = CE->getOperand(0);
> @@ -1004,8 +1005,8 @@ Constant *llvm::ConstantFoldInstOperands
> unsigned MidIntSize = CE->getType()->getScalarSizeInBits();
>
> if (MidIntSize >= SrcPtrSize) {
> - unsigned DestPtrSize = TD->getPointerTypeSizeInBits(DestTy);
> - if (SrcPtrSize == DestPtrSize)
> + unsigned SrcAS = SrcPtr->getType()->getPointerAddressSpace();
> + if (SrcAS == DestTy->getPointerAddressSpace())
> return FoldBitCast(CE->getOperand(0), DestTy, *TD);
> }
> }
> @@ -1021,6 +1022,7 @@ Constant *llvm::ConstantFoldInstOperands
> case Instruction::SIToFP:
> case Instruction::FPToUI:
> case Instruction::FPToSI:
> + case Instruction::AddrSpaceCast:
> return ConstantExpr::getCast(Opcode, Ops[0], DestTy);
> case Instruction::BitCast:
> if (TD)
>
> Modified: llvm/trunk/lib/AsmParser/LLLexer.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/AsmParser/LLLexer.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/AsmParser/LLLexer.cpp (original)
> +++ llvm/trunk/lib/AsmParser/LLLexer.cpp Thu Nov 14 19:34:59 2013
> @@ -665,6 +665,7 @@ lltok::Kind LLLexer::LexIdentifier() {
> INSTKEYWORD(inttoptr, IntToPtr);
> INSTKEYWORD(ptrtoint, PtrToInt);
> INSTKEYWORD(bitcast, BitCast);
> + INSTKEYWORD(addrspacecast, AddrSpaceCast);
> INSTKEYWORD(select, Select);
> INSTKEYWORD(va_arg, VAArg);
> INSTKEYWORD(ret, Ret);
>
> Modified: llvm/trunk/lib/AsmParser/LLParser.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/AsmParser/LLParser.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/AsmParser/LLParser.cpp (original)
> +++ llvm/trunk/lib/AsmParser/LLParser.cpp Thu Nov 14 19:34:59 2013
> @@ -2415,6 +2415,7 @@ bool LLParser::ParseValID(ValID &ID, Per
> case lltok::kw_fptrunc:
> case lltok::kw_fpext:
> case lltok::kw_bitcast:
> + case lltok::kw_addrspacecast:
> case lltok::kw_uitofp:
> case lltok::kw_sitofp:
> case lltok::kw_fptoui:
> @@ -3303,6 +3304,7 @@ int LLParser::ParseInstruction(Instructi
> case lltok::kw_fptrunc:
> case lltok::kw_fpext:
> case lltok::kw_bitcast:
> + case lltok::kw_addrspacecast:
> case lltok::kw_uitofp:
> case lltok::kw_sitofp:
> case lltok::kw_fptoui:
>
> Modified: llvm/trunk/lib/AsmParser/LLToken.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/AsmParser/LLToken.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/AsmParser/LLToken.h (original)
> +++ llvm/trunk/lib/AsmParser/LLToken.h Thu Nov 14 19:34:59 2013
> @@ -150,6 +150,7 @@ namespace lltok {
> kw_phi, kw_call,
> kw_trunc, kw_zext, kw_sext, kw_fptrunc, kw_fpext, kw_uitofp, kw_sitofp,
> kw_fptoui, kw_fptosi, kw_inttoptr, kw_ptrtoint, kw_bitcast,
> + kw_addrspacecast,
> kw_select, kw_va_arg,
>
> kw_landingpad, kw_personality, kw_cleanup, kw_catch, kw_filter,
>
> Modified: llvm/trunk/lib/Bitcode/Reader/BitcodeReader.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Bitcode/Reader/BitcodeReader.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Bitcode/Reader/BitcodeReader.cpp (original)
> +++ llvm/trunk/lib/Bitcode/Reader/BitcodeReader.cpp Thu Nov 14 19:34:59 2013
> @@ -128,6 +128,7 @@ static int GetDecodedCastOpcode(unsigned
> case bitc::CAST_PTRTOINT: return Instruction::PtrToInt;
> case bitc::CAST_INTTOPTR: return Instruction::IntToPtr;
> case bitc::CAST_BITCAST : return Instruction::BitCast;
> + case bitc::CAST_ADDRSPACECAST : return Instruction::AddrSpaceCast;
> }
> }
> static int GetDecodedBinaryOpcode(unsigned Val, Type *Ty) {
> @@ -1356,7 +1357,8 @@ error_code BitcodeReader::ParseConstants
> if (!OpTy)
> return Error(InvalidRecord);
> Constant *Op = ValueList.getConstantFwdRef(Record[2], OpTy);
> - V = ConstantExpr::getCast(Opc, Op, CurTy);
> + V = UpgradeBitCastExpr(Opc, Op, CurTy);
> + if (!V) V = ConstantExpr::getCast(Opc, Op, CurTy);
> }
> break;
> }
> @@ -2296,7 +2298,15 @@ error_code BitcodeReader::ParseFunctionB
> int Opc = GetDecodedCastOpcode(Record[OpNum+1]);
> if (Opc == -1 || ResTy == 0)
> return Error(InvalidRecord);
> - I = CastInst::Create((Instruction::CastOps)Opc, Op, ResTy);
> + Instruction *Temp = 0;
> + if ((I = UpgradeBitCastInst(Opc, Op, ResTy, Temp))) {
> + if (Temp) {
> + InstructionList.push_back(Temp);
> + CurBB->getInstList().push_back(Temp);
> + }
> + } else {
> + I = CastInst::Create((Instruction::CastOps)Opc, Op, ResTy);
> + }
> InstructionList.push_back(I);
> break;
> }
>
> Modified: llvm/trunk/lib/Bitcode/Writer/BitcodeWriter.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Bitcode/Writer/BitcodeWriter.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Bitcode/Writer/BitcodeWriter.cpp (original)
> +++ llvm/trunk/lib/Bitcode/Writer/BitcodeWriter.cpp Thu Nov 14 19:34:59 2013
> @@ -78,6 +78,7 @@ static unsigned GetEncodedCastOpcode(uns
> case Instruction::PtrToInt: return bitc::CAST_PTRTOINT;
> case Instruction::IntToPtr: return bitc::CAST_INTTOPTR;
> case Instruction::BitCast : return bitc::CAST_BITCAST;
> + case Instruction::AddrSpaceCast : return bitc::CAST_ADDRSPACECAST;
> }
> }
>
>
> Modified: llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAG.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAG.cpp (original)
> +++ llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAG.cpp Thu Nov 14 19:34:59 2013
> @@ -1507,6 +1507,26 @@ SDValue SelectionDAG::getMDNode(const MD
> return SDValue(N, 0);
> }
>
> +/// getAddrSpaceCast - Return an AddrSpaceCastSDNode.
> +SDValue SelectionDAG::getAddrSpaceCast(SDLoc dl, EVT VT, SDValue Ptr,
> + unsigned SrcAS, unsigned DestAS) {
> + SDValue Ops[] = {Ptr};
> + FoldingSetNodeID ID;
> + AddNodeIDNode(ID, ISD::ADDRSPACECAST, getVTList(VT), &Ops[0], 1);
> + ID.AddInteger(SrcAS);
> + ID.AddInteger(DestAS);
> +
> + void *IP = 0;
> + if (SDNode *E = CSEMap.FindNodeOrInsertPos(ID, IP))
> + return SDValue(E, 0);
> +
> + SDNode *N = new (NodeAllocator) AddrSpaceCastSDNode(dl.getIROrder(),
> + dl.getDebugLoc(),
> + VT, Ptr, SrcAS, DestAS);
> + CSEMap.InsertNode(N, IP);
> + AllNodes.push_back(N);
> + return SDValue(N, 0);
> +}
>
> /// getShiftAmountOperand - Return the specified value casted to
> /// the target's desired shift amount type.
> @@ -5978,6 +5998,12 @@ GlobalAddressSDNode::GlobalAddressSDNode
> TheGlobal = GA;
> }
>
> +AddrSpaceCastSDNode::AddrSpaceCastSDNode(unsigned Order, DebugLoc dl, EVT VT,
> + SDValue X, unsigned SrcAS,
> + unsigned DestAS)
> + : UnarySDNode(ISD::ADDRSPACECAST, Order, dl, getSDVTList(VT), X),
> + SrcAddrSpace(SrcAS), DestAddrSpace(DestAS) {}
> +
> MemSDNode::MemSDNode(unsigned Opc, unsigned Order, DebugLoc dl, SDVTList VTs,
> EVT memvt, MachineMemOperand *mmo)
> : SDNode(Opc, Order, dl, VTs), MemoryVT(memvt), MMO(mmo) {
>
> Modified: llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp (original)
> +++ llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp Thu Nov 14 19:34:59 2013
> @@ -2938,6 +2938,21 @@ void SelectionDAGBuilder::visitBitCast(c
> setValue(&I, N); // noop cast.
> }
>
> +void SelectionDAGBuilder::visitAddrSpaceCast(const User &I) {
> + const TargetLowering &TLI = DAG.getTargetLoweringInfo();
> + const Value *SV = I.getOperand(0);
> + SDValue N = getValue(SV);
> + EVT DestVT = TM.getTargetLowering()->getValueType(I.getType());
> +
> + unsigned SrcAS = SV->getType()->getPointerAddressSpace();
> + unsigned DestAS = I.getType()->getPointerAddressSpace();
> +
> + if (!TLI.isNoopAddrSpaceCast(SrcAS, DestAS))
> + N = DAG.getAddrSpaceCast(getCurSDLoc(), DestVT, N, SrcAS, DestAS);
> +
> + setValue(&I, N);
> +}
> +
> void SelectionDAGBuilder::visitInsertElement(const User &I) {
> const TargetLowering &TLI = DAG.getTargetLoweringInfo();
> SDValue InVec = getValue(I.getOperand(0));
>
> Modified: llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h (original)
> +++ llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h Thu Nov 14 19:34:59 2013
> @@ -26,6 +26,7 @@
>
> namespace llvm {
>
> +class AddrSpaceCastInst;
> class AliasAnalysis;
> class AllocaInst;
> class BasicBlock;
> @@ -720,6 +721,7 @@ private:
> void visitPtrToInt(const User &I);
> void visitIntToPtr(const User &I);
> void visitBitCast(const User &I);
> + void visitAddrSpaceCast(const User &I);
>
> void visitExtractElement(const User &I);
> void visitInsertElement(const User &I);
>
> Modified: llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp (original)
> +++ llvm/trunk/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp Thu Nov 14 19:34:59 2013
> @@ -224,6 +224,7 @@ std::string SDNode::getOperationName(con
> case ISD::FP_TO_SINT: return "fp_to_sint";
> case ISD::FP_TO_UINT: return "fp_to_uint";
> case ISD::BITCAST: return "bitcast";
> + case ISD::ADDRSPACECAST: return "addrspacecast";
> case ISD::FP16_TO_FP32: return "fp16_to_fp32";
> case ISD::FP32_TO_FP16: return "fp32_to_fp16";
>
> @@ -485,6 +486,13 @@ void SDNode::print_details(raw_ostream &
> OS << " " << offset;
> if (unsigned int TF = BA->getTargetFlags())
> OS << " [TF=" << TF << ']';
> + } else if (const AddrSpaceCastSDNode *ASC =
> + dyn_cast<AddrSpaceCastSDNode>(this)) {
> + OS << '['
> + << ASC->getSrcAddressSpace()
> + << " -> "
> + << ASC->getDestAddressSpace()
> + << ']';
> }
>
> if (unsigned Order = getIROrder())
>
> Modified: llvm/trunk/lib/CodeGen/TargetLoweringBase.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/TargetLoweringBase.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/TargetLoweringBase.cpp (original)
> +++ llvm/trunk/lib/CodeGen/TargetLoweringBase.cpp Thu Nov 14 19:34:59 2013
> @@ -1288,6 +1288,7 @@ int TargetLoweringBase::InstructionOpcod
> case PtrToInt: return ISD::BITCAST;
> case IntToPtr: return ISD::BITCAST;
> case BitCast: return ISD::BITCAST;
> + case AddrSpaceCast: return ISD::ADDRSPACECAST;
> case ICmp: return ISD::SETCC;
> case FCmp: return ISD::SETCC;
> case PHI: return 0;
>
> Modified: llvm/trunk/lib/IR/AutoUpgrade.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/AutoUpgrade.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/AutoUpgrade.cpp (original)
> +++ llvm/trunk/lib/IR/AutoUpgrade.cpp Thu Nov 14 19:34:59 2013
> @@ -447,3 +447,45 @@ void llvm::UpgradeInstWithTBAATag(Instru
> I->setMetadata(LLVMContext::MD_tbaa, MDNode::get(I->getContext(), Elts));
> }
> }
> +
> +Instruction *llvm::UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy,
> + Instruction *&Temp) {
> + if (Opc != Instruction::BitCast)
> + return 0;
> +
> + Temp = 0;
> + Type *SrcTy = V->getType();
> + if (SrcTy->isPtrOrPtrVectorTy() && DestTy->isPtrOrPtrVectorTy() &&
> + SrcTy->getPointerAddressSpace() != DestTy->getPointerAddressSpace()) {
> + LLVMContext &Context = V->getContext();
> +
> + // We have no information about target data layout, so we assume that
> + // the maximum pointer size is 64bit.
> + Type *MidTy = Type::getInt64Ty(Context);
> + Temp = CastInst::Create(Instruction::PtrToInt, V, MidTy);
> +
> + return CastInst::Create(Instruction::IntToPtr, Temp, DestTy);
> + }
> +
> + return 0;
> +}
> +
> +Value *llvm::UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy) {
> + if (Opc != Instruction::BitCast)
> + return 0;
> +
> + Type *SrcTy = C->getType();
> + if (SrcTy->isPtrOrPtrVectorTy() && DestTy->isPtrOrPtrVectorTy() &&
> + SrcTy->getPointerAddressSpace() != DestTy->getPointerAddressSpace()) {
> + LLVMContext &Context = C->getContext();
> +
> + // We have no information about target data layout, so we assume that
> + // the maximum pointer size is 64bit.
> + Type *MidTy = Type::getInt64Ty(Context);
> +
> + return ConstantExpr::getIntToPtr(ConstantExpr::getPtrToInt(C, MidTy),
> + DestTy);
> + }
> +
> + return 0;
> +}
>
> Modified: llvm/trunk/lib/IR/ConstantFold.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/ConstantFold.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/ConstantFold.cpp (original)
> +++ llvm/trunk/lib/IR/ConstantFold.cpp Thu Nov 14 19:34:59 2013
> @@ -689,6 +689,8 @@ Constant *llvm::ConstantFoldCastInstruct
> }
> case Instruction::BitCast:
> return FoldBitCast(V, DestTy);
> + case Instruction::AddrSpaceCast:
> + return 0;
> }
> }
>
>
> Modified: llvm/trunk/lib/IR/Constants.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Constants.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Constants.cpp (original)
> +++ llvm/trunk/lib/IR/Constants.cpp Thu Nov 14 19:34:59 2013
> @@ -1126,6 +1126,7 @@ getWithOperands(ArrayRef<Constant*> Ops,
> case Instruction::PtrToInt:
> case Instruction::IntToPtr:
> case Instruction::BitCast:
> + case Instruction::AddrSpaceCast:
> return ConstantExpr::getCast(getOpcode(), Ops[0], Ty);
> case Instruction::Select:
> return ConstantExpr::getSelect(Ops[0], Ops[1], Ops[2]);
> @@ -1461,6 +1462,7 @@ Constant *ConstantExpr::getCast(unsigned
> case Instruction::PtrToInt: return getPtrToInt(C, Ty);
> case Instruction::IntToPtr: return getIntToPtr(C, Ty);
> case Instruction::BitCast: return getBitCast(C, Ty);
> + case Instruction::AddrSpaceCast: return getAddrSpaceCast(C, Ty);
> }
> }
>
> @@ -1489,6 +1491,11 @@ Constant *ConstantExpr::getPointerCast(C
>
> if (Ty->isIntOrIntVectorTy())
> return getPtrToInt(S, Ty);
> +
> + unsigned SrcAS = S->getType()->getPointerAddressSpace();
> + if (Ty->isPtrOrPtrVectorTy() && SrcAS != Ty->getPointerAddressSpace())
> + return getAddrSpaceCast(S, Ty);
> +
> return getBitCast(S, Ty);
> }
>
> @@ -1662,6 +1669,13 @@ Constant *ConstantExpr::getBitCast(Const
> return getFoldedCast(Instruction::BitCast, C, DstTy);
> }
>
> +Constant *ConstantExpr::getAddrSpaceCast(Constant *C, Type *DstTy) {
> + assert(CastInst::castIsValid(Instruction::AddrSpaceCast, C, DstTy) &&
> + "Invalid constantexpr addrspacecast!");
> +
> + return getFoldedCast(Instruction::AddrSpaceCast, C, DstTy);
> +}
> +
> Constant *ConstantExpr::get(unsigned Opcode, Constant *C1, Constant *C2,
> unsigned Flags) {
> // Check the operands for consistency first.
>
> Modified: llvm/trunk/lib/IR/Core.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Core.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Core.cpp (original)
> +++ llvm/trunk/lib/IR/Core.cpp Thu Nov 14 19:34:59 2013
> @@ -1033,6 +1033,12 @@ LLVMValueRef LLVMConstBitCast(LLVMValueR
> unwrap(ToType)));
> }
>
> +LLVMValueRef LLVMConstAddrSpaceCast(LLVMValueRef ConstantVal,
> + LLVMTypeRef ToType) {
> + return wrap(ConstantExpr::getAddrSpaceCast(unwrap<Constant>(ConstantVal),
> + unwrap(ToType)));
> +}
> +
> LLVMValueRef LLVMConstZExtOrBitCast(LLVMValueRef ConstantVal,
> LLVMTypeRef ToType) {
> return wrap(ConstantExpr::getZExtOrBitCast(unwrap<Constant>(ConstantVal),
> @@ -2318,6 +2324,11 @@ LLVMValueRef LLVMBuildBitCast(LLVMBuilde
> return wrap(unwrap(B)->CreateBitCast(unwrap(Val), unwrap(DestTy), Name));
> }
>
> +LLVMValueRef LLVMBuildAddrSpaceCast(LLVMBuilderRef B, LLVMValueRef Val,
> + LLVMTypeRef DestTy, const char *Name) {
> + return wrap(unwrap(B)->CreateAddrSpaceCast(unwrap(Val), unwrap(DestTy), Name));
> +}
> +
> LLVMValueRef LLVMBuildZExtOrBitCast(LLVMBuilderRef B, LLVMValueRef Val,
> LLVMTypeRef DestTy, const char *Name) {
> return wrap(unwrap(B)->CreateZExtOrBitCast(unwrap(Val), unwrap(DestTy),
>
> Modified: llvm/trunk/lib/IR/Instruction.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Instruction.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Instruction.cpp (original)
> +++ llvm/trunk/lib/IR/Instruction.cpp Thu Nov 14 19:34:59 2013
> @@ -223,18 +223,19 @@ const char *Instruction::getOpcodeName(u
> case GetElementPtr: return "getelementptr";
>
> // Convert instructions...
> - case Trunc: return "trunc";
> - case ZExt: return "zext";
> - case SExt: return "sext";
> - case FPTrunc: return "fptrunc";
> - case FPExt: return "fpext";
> - case FPToUI: return "fptoui";
> - case FPToSI: return "fptosi";
> - case UIToFP: return "uitofp";
> - case SIToFP: return "sitofp";
> - case IntToPtr: return "inttoptr";
> - case PtrToInt: return "ptrtoint";
> - case BitCast: return "bitcast";
> + case Trunc: return "trunc";
> + case ZExt: return "zext";
> + case SExt: return "sext";
> + case FPTrunc: return "fptrunc";
> + case FPExt: return "fpext";
> + case FPToUI: return "fptoui";
> + case FPToSI: return "fptosi";
> + case UIToFP: return "uitofp";
> + case SIToFP: return "sitofp";
> + case IntToPtr: return "inttoptr";
> + case PtrToInt: return "ptrtoint";
> + case BitCast: return "bitcast";
> + case AddrSpaceCast: return "addrspacecast";
>
> // Other instructions...
> case ICmp: return "icmp";
>
> Modified: llvm/trunk/lib/IR/Instructions.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Instructions.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Instructions.cpp (original)
> +++ llvm/trunk/lib/IR/Instructions.cpp Thu Nov 14 19:34:59 2013
> @@ -2095,7 +2095,9 @@ bool CastInst::isNoopCast(Instruction::C
> case Instruction::SIToFP:
> case Instruction::FPToUI:
> case Instruction::FPToSI:
> - return false; // These always modify bits
> + case Instruction::AddrSpaceCast:
> + // TODO: Target informations may give a more accurate answer here.
> + return false;
> case Instruction::BitCast:
> return true; // BitCast never modifies bits.
> case Instruction::PtrToInt:
> @@ -2137,44 +2139,46 @@ unsigned CastInst::isEliminableCastPair(
> // ZEXT < Integral Unsigned Integer Any
> // SEXT < Integral Signed Integer Any
> // FPTOUI n/a FloatPt n/a Integral Unsigned
> - // FPTOSI n/a FloatPt n/a Integral Signed
> - // UITOFP n/a Integral Unsigned FloatPt n/a
> - // SITOFP n/a Integral Signed FloatPt n/a
> - // FPTRUNC > FloatPt n/a FloatPt n/a
> - // FPEXT < FloatPt n/a FloatPt n/a
> + // FPTOSI n/a FloatPt n/a Integral Signed
> + // UITOFP n/a Integral Unsigned FloatPt n/a
> + // SITOFP n/a Integral Signed FloatPt n/a
> + // FPTRUNC > FloatPt n/a FloatPt n/a
> + // FPEXT < FloatPt n/a FloatPt n/a
> // PTRTOINT n/a Pointer n/a Integral Unsigned
> // INTTOPTR n/a Integral Unsigned Pointer n/a
> - // BITCAST = FirstClass n/a FirstClass n/a
> + // BITCAST = FirstClass n/a FirstClass n/a
> + // ADDRSPCST n/a Pointer n/a Pointer n/a
> //
> // NOTE: some transforms are safe, but we consider them to be non-profitable.
> // For example, we could merge "fptoui double to i32" + "zext i32 to i64",
> // into "fptoui double to i64", but this loses information about the range
> - // of the produced value (we no longer know the top-part is all zeros).
> + // of the produced value (we no longer know the top-part is all zeros).
> // Further this conversion is often much more expensive for typical hardware,
> - // and causes issues when building libgcc. We disallow fptosi+sext for the
> + // and causes issues when building libgcc. We disallow fptosi+sext for the
> // same reason.
> - const unsigned numCastOps =
> + const unsigned numCastOps =
> Instruction::CastOpsEnd - Instruction::CastOpsBegin;
> static const uint8_t CastResults[numCastOps][numCastOps] = {
> - // T F F U S F F P I B -+
> - // R Z S P P I I T P 2 N T |
> - // U E E 2 2 2 2 R E I T C +- secondOp
> - // N X X U S F F N X N 2 V |
> - // C T T I I P P C T T P T -+
> - { 1, 0, 0,99,99, 0, 0,99,99,99, 0, 3 }, // Trunc -+
> - { 8, 1, 9,99,99, 2, 0,99,99,99, 2, 3 }, // ZExt |
> - { 8, 0, 1,99,99, 0, 2,99,99,99, 0, 3 }, // SExt |
> - { 0, 0, 0,99,99, 0, 0,99,99,99, 0, 3 }, // FPToUI |
> - { 0, 0, 0,99,99, 0, 0,99,99,99, 0, 3 }, // FPToSI |
> - { 99,99,99, 0, 0,99,99, 0, 0,99,99, 4 }, // UIToFP +- firstOp
> - { 99,99,99, 0, 0,99,99, 0, 0,99,99, 4 }, // SIToFP |
> - { 99,99,99, 0, 0,99,99, 1, 0,99,99, 4 }, // FPTrunc |
> - { 99,99,99, 2, 2,99,99,10, 2,99,99, 4 }, // FPExt |
> - { 1, 0, 0,99,99, 0, 0,99,99,99, 7, 3 }, // PtrToInt |
> - { 99,99,99,99,99,99,99,99,99,13,99,12 }, // IntToPtr |
> - { 5, 5, 5, 6, 6, 5, 5, 6, 6,11, 5, 1 }, // BitCast -+
> + // T F F U S F F P I B A -+
> + // R Z S P P I I T P 2 N T S |
> + // U E E 2 2 2 2 R E I T C C +- secondOp
> + // N X X U S F F N X N 2 V V |
> + // C T T I I P P C T T P T T -+
> + { 1, 0, 0,99,99, 0, 0,99,99,99, 0, 3, 0}, // Trunc -+
> + { 8, 1, 9,99,99, 2, 0,99,99,99, 2, 3, 0}, // ZExt |
> + { 8, 0, 1,99,99, 0, 2,99,99,99, 0, 3, 0}, // SExt |
> + { 0, 0, 0,99,99, 0, 0,99,99,99, 0, 3, 0}, // FPToUI |
> + { 0, 0, 0,99,99, 0, 0,99,99,99, 0, 3, 0}, // FPToSI |
> + { 99,99,99, 0, 0,99,99, 0, 0,99,99, 4, 0}, // UIToFP +- firstOp
> + { 99,99,99, 0, 0,99,99, 0, 0,99,99, 4, 0}, // SIToFP |
> + { 99,99,99, 0, 0,99,99, 1, 0,99,99, 4, 0}, // FPTrunc |
> + { 99,99,99, 2, 2,99,99,10, 2,99,99, 4, 0}, // FPExt |
> + { 1, 0, 0,99,99, 0, 0,99,99,99, 7, 3, 0}, // PtrToInt |
> + { 99,99,99,99,99,99,99,99,99,11,99,15, 0}, // IntToPtr |
> + { 5, 5, 5, 6, 6, 5, 5, 6, 6,16, 5, 1,14}, // BitCast |
> + { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,13,12}, // AddrSpaceCast -+
> };
> -
> +
> // If either of the casts are a bitcast from scalar to vector, disallow the
> // merging. However, bitcast of A->B->A are allowed.
> bool isFirstBitcast = (firstOp == Instruction::BitCast);
> @@ -2191,47 +2195,50 @@ unsigned CastInst::isEliminableCastPair(
> [secondOp-Instruction::CastOpsBegin];
> switch (ElimCase) {
> case 0:
> - // categorically disallowed
> + // Categorically disallowed.
> return 0;
> case 1:
> - // allowed, use first cast's opcode
> + // Allowed, use first cast's opcode.
> return firstOp;
> case 2:
> - // allowed, use second cast's opcode
> + // Allowed, use second cast's opcode.
> return secondOp;
> case 3:
> - // no-op cast in second op implies firstOp as long as the DestTy
> + // No-op cast in second op implies firstOp as long as the DestTy
> // is integer and we are not converting between a vector and a
> // non vector type.
> if (!SrcTy->isVectorTy() && DstTy->isIntegerTy())
> return firstOp;
> return 0;
> case 4:
> - // no-op cast in second op implies firstOp as long as the DestTy
> + // No-op cast in second op implies firstOp as long as the DestTy
> // is floating point.
> if (DstTy->isFloatingPointTy())
> return firstOp;
> return 0;
> case 5:
> - // no-op cast in first op implies secondOp as long as the SrcTy
> + // No-op cast in first op implies secondOp as long as the SrcTy
> // is an integer.
> if (SrcTy->isIntegerTy())
> return secondOp;
> return 0;
> case 6:
> - // no-op cast in first op implies secondOp as long as the SrcTy
> + // No-op cast in first op implies secondOp as long as the SrcTy
> // is a floating point.
> if (SrcTy->isFloatingPointTy())
> return secondOp;
> return 0;
> case 7: {
> + // Cannot simplify if address spaces are different!
> + if (SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace())
> + return 0;
> +
> unsigned MidSize = MidTy->getScalarSizeInBits();
> - // Check the address spaces first. If we know they are in the same address
> - // space, the pointer sizes must be the same so we can still fold this
> - // without knowing the actual sizes as long we know that the intermediate
> - // pointer is the largest possible pointer size.
> - if (MidSize == 64 &&
> - SrcTy->getPointerAddressSpace() == DstTy->getPointerAddressSpace())
> + // We can still fold this without knowing the actual sizes as long we
> + // know that the intermediate pointer is the largest possible
> + // pointer size.
> + // FIXME: Is this always true?
> + if (MidSize == 64)
> return Instruction::BitCast;
>
> // ptrtoint, inttoptr -> bitcast (ptr -> ptr) if int size is >= ptr size.
> @@ -2254,7 +2261,8 @@ unsigned CastInst::isEliminableCastPair(
> return firstOp;
> return secondOp;
> }
> - case 9: // zext, sext -> zext, because sext can't sign extend after zext
> + case 9:
> + // zext, sext -> zext, because sext can't sign extend after zext
> return Instruction::ZExt;
> case 10:
> // fpext followed by ftrunc is allowed if the bit size returned to is
> @@ -2263,46 +2271,6 @@ unsigned CastInst::isEliminableCastPair(
> return Instruction::BitCast;
> return 0; // If the types are not the same we can't eliminate it.
> case 11: {
> - // bitcast followed by ptrtoint is allowed as long as the bitcast is a
> - // pointer to pointer cast, and the pointers are the same size.
> - PointerType *SrcPtrTy = dyn_cast<PointerType>(SrcTy);
> - PointerType *MidPtrTy = dyn_cast<PointerType>(MidTy);
> - if (!SrcPtrTy || !MidPtrTy)
> - return 0;
> -
> - // If the address spaces are the same, we know they are the same size
> - // without size information
> - if (SrcPtrTy->getAddressSpace() == MidPtrTy->getAddressSpace())
> - return secondOp;
> -
> - if (!SrcIntPtrTy || !MidIntPtrTy)
> - return 0;
> -
> - if (SrcIntPtrTy->getScalarSizeInBits() ==
> - MidIntPtrTy->getScalarSizeInBits())
> - return secondOp;
> -
> - return 0;
> - }
> - case 12: {
> - // inttoptr, bitcast -> inttoptr if bitcast is a ptr to ptr cast
> - // and the ptrs are to address spaces of the same size
> - PointerType *MidPtrTy = dyn_cast<PointerType>(MidTy);
> - PointerType *DstPtrTy = dyn_cast<PointerType>(DstTy);
> - if (!MidPtrTy || !DstPtrTy)
> - return 0;
> -
> - if (MidPtrTy->getAddressSpace() == DstPtrTy->getAddressSpace())
> - return firstOp;
> -
> - if (MidIntPtrTy &&
> - DstIntPtrTy &&
> - MidIntPtrTy->getScalarSizeInBits() ==
> - DstIntPtrTy->getScalarSizeInBits())
> - return firstOp;
> - return 0;
> - }
> - case 13: {
> // inttoptr, ptrtoint -> bitcast if SrcSize<=PtrSize and SrcSize==DstSize
> if (!MidIntPtrTy)
> return 0;
> @@ -2313,8 +2281,65 @@ unsigned CastInst::isEliminableCastPair(
> return Instruction::BitCast;
> return 0;
> }
> + case 12: {
> + // addrspacecast, addrspacecast -> bitcast, if SrcAS == DstAS
> + // addrspacecast, addrspacecast -> addrspacecast, if SrcAS != DstAS
> + if (SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace())
> + return Instruction::AddrSpaceCast;
> + return Instruction::BitCast;
> + }
> + case 13:
> + // FIXME: this state can be merged with (1), but the following assert
> + // is useful to check the correcteness of the sequence due to semantic
> + // change of bitcast.
> + assert(
> + SrcTy->isPtrOrPtrVectorTy() &&
> + MidTy->isPtrOrPtrVectorTy() &&
> + DstTy->isPtrOrPtrVectorTy() &&
> + SrcTy->getPointerAddressSpace() != MidTy->getPointerAddressSpace() &&
> + MidTy->getPointerAddressSpace() == DstTy->getPointerAddressSpace() &&
> + "Illegal addrspacecast, bitcast sequence!");
> + // Allowed, use first cast's opcode
> + return firstOp;
> + case 14:
> + // FIXME: this state can be merged with (2), but the following assert
> + // is useful to check the correcteness of the sequence due to semantic
> + // change of bitcast.
> + assert(
> + SrcTy->isPtrOrPtrVectorTy() &&
> + MidTy->isPtrOrPtrVectorTy() &&
> + DstTy->isPtrOrPtrVectorTy() &&
> + SrcTy->getPointerAddressSpace() == MidTy->getPointerAddressSpace() &&
> + MidTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace() &&
> + "Illegal bitcast, addrspacecast sequence!");
> + // Allowed, use second cast's opcode
> + return secondOp;
> + case 15:
> + // FIXME: this state can be merged with (1), but the following assert
> + // is useful to check the correcteness of the sequence due to semantic
> + // change of bitcast.
> + assert(
> + SrcTy->isIntOrIntVectorTy() &&
> + MidTy->isPtrOrPtrVectorTy() &&
> + DstTy->isPtrOrPtrVectorTy() &&
> + MidTy->getPointerAddressSpace() == DstTy->getPointerAddressSpace() &&
> + "Illegal inttoptr, bitcast sequence!");
> + // Allowed, use first cast's opcode
> + return firstOp;
> + case 16:
> + // FIXME: this state can be merged with (2), but the following assert
> + // is useful to check the correcteness of the sequence due to semantic
> + // change of bitcast.
> + assert(
> + SrcTy->isPtrOrPtrVectorTy() &&
> + MidTy->isPtrOrPtrVectorTy() &&
> + DstTy->isIntOrIntVectorTy() &&
> + SrcTy->getPointerAddressSpace() == MidTy->getPointerAddressSpace() &&
> + "Illegal bitcast, ptrtoint sequence!");
> + // Allowed, use second cast's opcode
> + return secondOp;
> case 99:
> - // cast combination can't happen (error in input). This is for all cases
> + // Cast combination can't happen (error in input). This is for all cases
> // where the MidTy is not the same for the two cast instructions.
> llvm_unreachable("Invalid Cast Combination");
> default:
> @@ -2327,19 +2352,20 @@ CastInst *CastInst::Create(Instruction::
> assert(castIsValid(op, S, Ty) && "Invalid cast!");
> // Construct and return the appropriate CastInst subclass
> switch (op) {
> - case Trunc: return new TruncInst (S, Ty, Name, InsertBefore);
> - case ZExt: return new ZExtInst (S, Ty, Name, InsertBefore);
> - case SExt: return new SExtInst (S, Ty, Name, InsertBefore);
> - case FPTrunc: return new FPTruncInst (S, Ty, Name, InsertBefore);
> - case FPExt: return new FPExtInst (S, Ty, Name, InsertBefore);
> - case UIToFP: return new UIToFPInst (S, Ty, Name, InsertBefore);
> - case SIToFP: return new SIToFPInst (S, Ty, Name, InsertBefore);
> - case FPToUI: return new FPToUIInst (S, Ty, Name, InsertBefore);
> - case FPToSI: return new FPToSIInst (S, Ty, Name, InsertBefore);
> - case PtrToInt: return new PtrToIntInst (S, Ty, Name, InsertBefore);
> - case IntToPtr: return new IntToPtrInst (S, Ty, Name, InsertBefore);
> - case BitCast: return new BitCastInst (S, Ty, Name, InsertBefore);
> - default: llvm_unreachable("Invalid opcode provided");
> + case Trunc: return new TruncInst (S, Ty, Name, InsertBefore);
> + case ZExt: return new ZExtInst (S, Ty, Name, InsertBefore);
> + case SExt: return new SExtInst (S, Ty, Name, InsertBefore);
> + case FPTrunc: return new FPTruncInst (S, Ty, Name, InsertBefore);
> + case FPExt: return new FPExtInst (S, Ty, Name, InsertBefore);
> + case UIToFP: return new UIToFPInst (S, Ty, Name, InsertBefore);
> + case SIToFP: return new SIToFPInst (S, Ty, Name, InsertBefore);
> + case FPToUI: return new FPToUIInst (S, Ty, Name, InsertBefore);
> + case FPToSI: return new FPToSIInst (S, Ty, Name, InsertBefore);
> + case PtrToInt: return new PtrToIntInst (S, Ty, Name, InsertBefore);
> + case IntToPtr: return new IntToPtrInst (S, Ty, Name, InsertBefore);
> + case BitCast: return new BitCastInst (S, Ty, Name, InsertBefore);
> + case AddrSpaceCast: return new AddrSpaceCastInst (S, Ty, Name, InsertBefore);
> + default: llvm_unreachable("Invalid opcode provided");
> }
> }
>
> @@ -2348,19 +2374,20 @@ CastInst *CastInst::Create(Instruction::
> assert(castIsValid(op, S, Ty) && "Invalid cast!");
> // Construct and return the appropriate CastInst subclass
> switch (op) {
> - case Trunc: return new TruncInst (S, Ty, Name, InsertAtEnd);
> - case ZExt: return new ZExtInst (S, Ty, Name, InsertAtEnd);
> - case SExt: return new SExtInst (S, Ty, Name, InsertAtEnd);
> - case FPTrunc: return new FPTruncInst (S, Ty, Name, InsertAtEnd);
> - case FPExt: return new FPExtInst (S, Ty, Name, InsertAtEnd);
> - case UIToFP: return new UIToFPInst (S, Ty, Name, InsertAtEnd);
> - case SIToFP: return new SIToFPInst (S, Ty, Name, InsertAtEnd);
> - case FPToUI: return new FPToUIInst (S, Ty, Name, InsertAtEnd);
> - case FPToSI: return new FPToSIInst (S, Ty, Name, InsertAtEnd);
> - case PtrToInt: return new PtrToIntInst (S, Ty, Name, InsertAtEnd);
> - case IntToPtr: return new IntToPtrInst (S, Ty, Name, InsertAtEnd);
> - case BitCast: return new BitCastInst (S, Ty, Name, InsertAtEnd);
> - default: llvm_unreachable("Invalid opcode provided");
> + case Trunc: return new TruncInst (S, Ty, Name, InsertAtEnd);
> + case ZExt: return new ZExtInst (S, Ty, Name, InsertAtEnd);
> + case SExt: return new SExtInst (S, Ty, Name, InsertAtEnd);
> + case FPTrunc: return new FPTruncInst (S, Ty, Name, InsertAtEnd);
> + case FPExt: return new FPExtInst (S, Ty, Name, InsertAtEnd);
> + case UIToFP: return new UIToFPInst (S, Ty, Name, InsertAtEnd);
> + case SIToFP: return new SIToFPInst (S, Ty, Name, InsertAtEnd);
> + case FPToUI: return new FPToUIInst (S, Ty, Name, InsertAtEnd);
> + case FPToSI: return new FPToSIInst (S, Ty, Name, InsertAtEnd);
> + case PtrToInt: return new PtrToIntInst (S, Ty, Name, InsertAtEnd);
> + case IntToPtr: return new IntToPtrInst (S, Ty, Name, InsertAtEnd);
> + case BitCast: return new BitCastInst (S, Ty, Name, InsertAtEnd);
> + case AddrSpaceCast: return new AddrSpaceCastInst (S, Ty, Name, InsertAtEnd);
> + default: llvm_unreachable("Invalid opcode provided");
> }
> }
>
> @@ -2425,6 +2452,11 @@ CastInst *CastInst::CreatePointerCast(Va
>
> if (Ty->isIntOrIntVectorTy())
> return Create(Instruction::PtrToInt, S, Ty, Name, InsertAtEnd);
> +
> + Type *STy = S->getType();
> + if (STy->getPointerAddressSpace() != Ty->getPointerAddressSpace())
> + return Create(Instruction::AddrSpaceCast, S, Ty, Name, InsertAtEnd);
> +
> return Create(Instruction::BitCast, S, Ty, Name, InsertAtEnd);
> }
>
> @@ -2442,6 +2474,11 @@ CastInst *CastInst::CreatePointerCast(Va
>
> if (Ty->isIntOrIntVectorTy())
> return Create(Instruction::PtrToInt, S, Ty, Name, InsertBefore);
> +
> + Type *STy = S->getType();
> + if (STy->getPointerAddressSpace() != Ty->getPointerAddressSpace())
> + return Create(Instruction::AddrSpaceCast, S, Ty, Name, InsertBefore);
> +
> return Create(Instruction::BitCast, S, Ty, Name, InsertBefore);
> }
>
> @@ -2687,7 +2724,8 @@ CastInst::getCastOpcode(
> return BitCast;
> } else if (DestTy->isPointerTy()) {
> if (SrcTy->isPointerTy()) {
> - // TODO: Address space pointer sizes may not match
> + if (DestTy->getPointerAddressSpace() != SrcTy->getPointerAddressSpace())
> + return AddrSpaceCast;
> return BitCast; // ptr -> ptr
> } else if (SrcTy->isIntegerTy()) {
> return IntToPtr; // int -> ptr
> @@ -2782,13 +2820,27 @@ CastInst::castIsValid(Instruction::CastO
> case Instruction::BitCast:
> // BitCast implies a no-op cast of type only. No bits change.
> // However, you can't cast pointers to anything but pointers.
> - if (SrcTy->isPointerTy() != DstTy->isPointerTy())
> + if (SrcTy->isPtrOrPtrVectorTy() != DstTy->isPtrOrPtrVectorTy())
> return false;
>
> - // Now we know we're not dealing with a pointer/non-pointer mismatch. In all
> - // these cases, the cast is okay if the source and destination bit widths
> - // are identical.
> - return SrcTy->getPrimitiveSizeInBits() == DstTy->getPrimitiveSizeInBits();
> + // For non pointer cases, the cast is okay if the source and destination bit
> + // widths are identical.
> + if (!SrcTy->isPtrOrPtrVectorTy())
> + return SrcTy->getPrimitiveSizeInBits() == DstTy->getPrimitiveSizeInBits();
> +
> + // If both are pointers then the address spaces must match and vector of
> + // pointers must have the same number of elements.
> + return SrcTy->getPointerAddressSpace() == DstTy->getPointerAddressSpace() &&
> + SrcTy->isVectorTy() == DstTy->isVectorTy() &&
> + (!SrcTy->isVectorTy() ||
> + SrcTy->getVectorNumElements() == SrcTy->getVectorNumElements());
> +
> + case Instruction::AddrSpaceCast:
> + return SrcTy->isPtrOrPtrVectorTy() && DstTy->isPtrOrPtrVectorTy() &&
> + SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace() &&
> + SrcTy->isVectorTy() == DstTy->isVectorTy() &&
> + (!SrcTy->isVectorTy() ||
> + SrcTy->getVectorNumElements() == SrcTy->getVectorNumElements());
> }
> }
>
> @@ -2935,6 +2987,18 @@ BitCastInst::BitCastInst(
> assert(castIsValid(getOpcode(), S, Ty) && "Illegal BitCast");
> }
>
> +AddrSpaceCastInst::AddrSpaceCastInst(
> + Value *S, Type *Ty, const Twine &Name, Instruction *InsertBefore
> +) : CastInst(Ty, AddrSpaceCast, S, Name, InsertBefore) {
> + assert(castIsValid(getOpcode(), S, Ty) && "Illegal AddrSpaceCast");
> +}
> +
> +AddrSpaceCastInst::AddrSpaceCastInst(
> + Value *S, Type *Ty, const Twine &Name, BasicBlock *InsertAtEnd
> +) : CastInst(Ty, AddrSpaceCast, S, Name, InsertAtEnd) {
> + assert(castIsValid(getOpcode(), S, Ty) && "Illegal AddrSpaceCast");
> +}
> +
> //===----------------------------------------------------------------------===//
> // CmpInst Classes
> //===----------------------------------------------------------------------===//
> @@ -3553,6 +3617,10 @@ BitCastInst *BitCastInst::clone_impl() c
> return new BitCastInst(getOperand(0), getType());
> }
>
> +AddrSpaceCastInst *AddrSpaceCastInst::clone_impl() const {
> + return new AddrSpaceCastInst(getOperand(0), getType());
> +}
> +
> CallInst *CallInst::clone_impl() const {
> return new(getNumOperands()) CallInst(*this);
> }
>
> Modified: llvm/trunk/lib/IR/Value.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Value.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Value.cpp (original)
> +++ llvm/trunk/lib/IR/Value.cpp Thu Nov 14 19:34:59 2013
> @@ -365,7 +365,8 @@ static Value *stripPointerCastsAndOffset
> break;
> }
> V = GEP->getPointerOperand();
> - } else if (Operator::getOpcode(V) == Instruction::BitCast) {
> + } else if (Operator::getOpcode(V) == Instruction::BitCast ||
> + Operator::getOpcode(V) == Instruction::AddrSpaceCast) {
> V = cast<Operator>(V)->getOperand(0);
> } else if (GlobalAlias *GA = dyn_cast<GlobalAlias>(V)) {
> if (StripKind == PSK_ZeroIndices || GA->mayBeOverridden())
>
> Modified: llvm/trunk/lib/IR/Verifier.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/IR/Verifier.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/IR/Verifier.cpp (original)
> +++ llvm/trunk/lib/IR/Verifier.cpp Thu Nov 14 19:34:59 2013
> @@ -283,6 +283,7 @@ namespace {
> void visitIntToPtrInst(IntToPtrInst &I);
> void visitPtrToIntInst(PtrToIntInst &I);
> void visitBitCastInst(BitCastInst &I);
> + void visitAddrSpaceCastInst(AddrSpaceCastInst &I);
> void visitPHINode(PHINode &PN);
> void visitBinaryOperator(BinaryOperator &B);
> void visitICmpInst(ICmpInst &IC);
> @@ -965,11 +966,9 @@ void Verifier::VerifyBitcastType(const V
> unsigned SrcAS = SrcTy->getPointerAddressSpace();
> unsigned DstAS = DestTy->getPointerAddressSpace();
>
> - unsigned SrcASSize = DL->getPointerSizeInBits(SrcAS);
> - unsigned DstASSize = DL->getPointerSizeInBits(DstAS);
> - Assert1(SrcASSize == DstASSize,
> - "Bitcasts between pointers of different address spaces must have "
> - "the same size pointers, otherwise use PtrToInt/IntToPtr.", V);
> + Assert1(SrcAS == DstAS,
> + "Bitcasts between pointers of different address spaces is not legal."
> + "Use AddrSpaceCast instead.", V);
> }
>
> void Verifier::VerifyConstantExprBitcastType(const ConstantExpr *CE) {
> @@ -1455,6 +1454,22 @@ void Verifier::visitBitCastInst(BitCastI
> visitInstruction(I);
> }
>
> +void Verifier::visitAddrSpaceCastInst(AddrSpaceCastInst &I) {
> + Type *SrcTy = I.getOperand(0)->getType();
> + Type *DestTy = I.getType();
> +
> + Assert1(SrcTy->isPtrOrPtrVectorTy(),
> + "AddrSpaceCast source must be a pointer", &I);
> + Assert1(DestTy->isPtrOrPtrVectorTy(),
> + "AddrSpaceCast result must be a pointer", &I);
> + Assert1(SrcTy->getPointerAddressSpace() != DestTy->getPointerAddressSpace(),
> + "AddrSpaceCast must be between different address spaces", &I);
> + if (SrcTy->isVectorTy())
> + Assert1(SrcTy->getVectorNumElements() == DestTy->getVectorNumElements(),
> + "AddrSpaceCast vector pointer number of elements mismatch", &I);
> + visitInstruction(I);
> +}
> +
> /// visitPHINode - Ensure that a PHI node is well formed.
> ///
> void Verifier::visitPHINode(PHINode &PN) {
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp Thu Nov 14 19:34:59 2013
> @@ -142,7 +142,7 @@ bool GenericToNVVM::runOnModule(Module &
> GlobalVariable *GV = I->first;
> GlobalVariable *NewGV = I->second;
> ++I;
> - Constant *BitCastNewGV = ConstantExpr::getBitCast(NewGV, GV->getType());
> + Constant *BitCastNewGV = ConstantExpr::getPointerCast(NewGV, GV->getType());
> // At this point, the remaining uses of GV should be found only in global
> // variable initializers, as other uses have been already been removed
> // while walking through the instructions in function definitions.
>
> Modified: llvm/trunk/lib/Target/X86/X86ISelLowering.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/X86/X86ISelLowering.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Target/X86/X86ISelLowering.cpp (original)
> +++ llvm/trunk/lib/Target/X86/X86ISelLowering.cpp Thu Nov 14 19:34:59 2013
> @@ -92,7 +92,7 @@ static SDValue ExtractSubVector(SDValue
> VecIdx);
>
> return Result;
> -
> +
> }
> /// Generate a DAG to grab 128-bits from a vector > 128 bits. This
> /// sets things up to match to an AVX VEXTRACTF128 / VEXTRACTI128
> @@ -1765,6 +1765,13 @@ bool X86TargetLowering::getStackCookieLo
> return true;
> }
>
> +bool X86TargetLowering::isNoopAddrSpaceCast(unsigned SrcAS,
> + unsigned DestAS) const {
> + assert(SrcAS != DestAS && "Expected different address spaces!");
> +
> + return SrcAS < 256 && DestAS < 256;
> +}
> +
> //===----------------------------------------------------------------------===//
> // Return Value Calling Convention Implementation
> //===----------------------------------------------------------------------===//
>
> Modified: llvm/trunk/lib/Target/X86/X86ISelLowering.h
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/X86/X86ISelLowering.h?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/lib/Target/X86/X86ISelLowering.h (original)
> +++ llvm/trunk/lib/Target/X86/X86ISelLowering.h Thu Nov 14 19:34:59 2013
> @@ -775,6 +775,8 @@ namespace llvm {
> SDValue BuildFILD(SDValue Op, EVT SrcVT, SDValue Chain, SDValue StackSlot,
> SelectionDAG &DAG) const;
>
> + virtual bool isNoopAddrSpaceCast(unsigned SrcAS, unsigned DestAS) const LLVM_OVERRIDE;
> +
> /// \brief Reset the operation actions based on target options.
> virtual void resetOperationActions();
>
>
> Modified: llvm/trunk/test/Assembler/ConstantExprFoldCast.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Assembler/ConstantExprFoldCast.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Assembler/ConstantExprFoldCast.ll (original)
> +++ llvm/trunk/test/Assembler/ConstantExprFoldCast.ll Thu Nov 14 19:34:59 2013
> @@ -12,3 +12,5 @@
> @F = global i32* inttoptr (i32 add (i32 5, i32 -5) to i32*)
> @G = global i32* inttoptr (i32 sub (i32 5, i32 5) to i32*)
>
> +; Address space cast AS0 null-> AS1 null
> + at H = global i32 addrspace(1)* addrspacecast(i32* null to i32 addrspace(1)*)
>
> Modified: llvm/trunk/test/Assembler/ConstantExprNoFold.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Assembler/ConstantExprNoFold.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Assembler/ConstantExprNoFold.ll (original)
> +++ llvm/trunk/test/Assembler/ConstantExprNoFold.ll Thu Nov 14 19:34:59 2013
> @@ -21,3 +21,6 @@ target datalayout = "p:32:32"
>
> ; CHECK: @D = global i1 icmp eq (i64* getelementptr inbounds (i64* @A, i64 1), i64* getelementptr inbounds (i64* @B, i64 2))
> @D = global i1 icmp eq (i64* getelementptr inbounds (i64* @A, i64 1), i64* getelementptr inbounds (i64* @B, i64 2))
> +
> +; CHECK: @E = global i64 addrspace(1)* addrspacecast (i64* @A to i64 addrspace(1)*)
> + at E = global i64 addrspace(1)* addrspacecast(i64* @A to i64 addrspace(1)*)
>
> Modified: llvm/trunk/test/CodeGen/X86/atomic-dagsched.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/X86/atomic-dagsched.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/CodeGen/X86/atomic-dagsched.ll (original)
> +++ llvm/trunk/test/CodeGen/X86/atomic-dagsched.ll Thu Nov 14 19:34:59 2013
> @@ -34,8 +34,8 @@ dim_0_vector_pre_head.i:
> vector_kernel_entry.i: ; preds = %vector_kernel_entry.i, %dim_0_vector_pre_head.i
> %asr.iv9 = phi i8* [ %scevgep10, %vector_kernel_entry.i ], [ %asr.iv6, %dim_0_vector_pre_head.i ]
> %asr.iv = phi i64 [ %asr.iv.next, %vector_kernel_entry.i ], [ %vector.size.i, %dim_0_vector_pre_head.i ]
> - %8 = bitcast i8* %ptrtoarg4 to i32 addrspace(1)*
> - %asr.iv911 = bitcast i8* %asr.iv9 to <8 x i32> addrspace(1)*
> + %8 = addrspacecast i8* %ptrtoarg4 to i32 addrspace(1)*
> + %asr.iv911 = addrspacecast i8* %asr.iv9 to <8 x i32> addrspace(1)*
> %9 = load <8 x i32> addrspace(1)* %asr.iv911, align 4
> %extract8vector_func.i = extractelement <8 x i32> %9, i32 0
> %extract9vector_func.i = extractelement <8 x i32> %9, i32 1
> @@ -73,8 +73,8 @@ dim_0_pre_head.i:
>
> scalar_kernel_entry.i: ; preds = %scalar_kernel_entry.i, %dim_0_pre_head.i
> %asr.iv12 = phi i64 [ %asr.iv.next13, %scalar_kernel_entry.i ], [ %22, %dim_0_pre_head.i ]
> - %23 = bitcast i8* %asr.iv6 to i32 addrspace(1)*
> - %24 = bitcast i8* %ptrtoarg4 to i32 addrspace(1)*
> + %23 = addrspacecast i8* %asr.iv6 to i32 addrspace(1)*
> + %24 = addrspacecast i8* %ptrtoarg4 to i32 addrspace(1)*
> %scevgep16 = getelementptr i32 addrspace(1)* %23, i64 %asr.iv12
> %25 = load i32 addrspace(1)* %scevgep16, align 4
> %26 = atomicrmw min i32 addrspace(1)* %24, i32 %25 seq_cst
>
> Modified: llvm/trunk/test/DebugInfo/X86/dbg-value-isel.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/X86/dbg-value-isel.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/DebugInfo/X86/dbg-value-isel.ll (original)
> +++ llvm/trunk/test/DebugInfo/X86/dbg-value-isel.ll Thu Nov 14 19:34:59 2013
> @@ -9,7 +9,7 @@ target triple = "x86_64-apple-darwin10.0
> @sgv = internal addrspace(2) constant [1 x i8] zeroinitializer
> @fgv = internal addrspace(2) constant [1 x i8] zeroinitializer
> @lvgv = internal constant [0 x i8*] zeroinitializer
> - at llvm.global.annotations = appending global [1 x %0] [%0 { i8* bitcast (void (i32 addrspace(1)*)* @__OpenCL_nbt02_kernel to i8*), i8* bitcast ([1 x i8] addrspace(2)* @sgv to i8*), i8* bitcast ([1 x i8] addrspace(2)* @fgv to i8*), i8* bitcast ([0 x i8*]* @lvgv to i8*), i32 0 }], section "llvm.metadata"
> + at llvm.global.annotations = appending global [1 x %0] [%0 { i8* bitcast (void (i32 addrspace(1)*)* @__OpenCL_nbt02_kernel to i8*), i8* addrspacecast ([1 x i8] addrspace(2)* @sgv to i8*), i8* addrspacecast ([1 x i8] addrspace(2)* @fgv to i8*), i8* bitcast ([0 x i8*]* @lvgv to i8*), i32 0 }], section "llvm.metadata"
>
> define void @__OpenCL_nbt02_kernel(i32 addrspace(1)* %ip) nounwind {
> entry:
>
> Modified: llvm/trunk/test/Feature/newcasts.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Feature/newcasts.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Feature/newcasts.ll (original)
> +++ llvm/trunk/test/Feature/newcasts.ll Thu Nov 14 19:34:59 2013
> @@ -20,6 +20,9 @@ define void @"NewCasts" (i16 %x) {
> %p = uitofp <4 x i32> %n to <4 x float>
> %q = fptosi <4 x float> %p to <4 x i32>
> %r = fptoui <4 x float> %p to <4 x i32>
> + %s = inttoptr <4 x i32> %n to <4 x i32*>
> + %t = addrspacecast <4 x i32*> %s to <4 x i32 addrspace(1)*>
> + %z = addrspacecast <4 x i32*> %s to <4 x float addrspace(2)*>
> ret void
> }
>
>
> 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=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Other/constant-fold-gep.ll (original)
> +++ llvm/trunk/test/Other/constant-fold-gep.ll Thu Nov 14 19:34:59 2013
> @@ -454,10 +454,10 @@ define i32* @fZ() nounwind {
>
> define i8* @different_addrspace() nounwind noinline {
> ; OPT: different_addrspace
> - %p = getelementptr inbounds i8* bitcast ([4 x i8] addrspace(12)* @p12 to i8*),
> + %p = getelementptr inbounds i8* addrspacecast ([4 x i8] addrspace(12)* @p12 to i8*),
> i32 2
> ret i8* %p
> -; OPT: ret i8* getelementptr (i8* bitcast ([4 x i8] addrspace(12)* @p12 to i8*), i32 2)
> +; OPT: ret i8* getelementptr (i8* addrspacecast ([4 x i8] addrspace(12)* @p12 to i8*), i32 2)
> }
>
> define i8* @same_addrspace() nounwind noinline {
>
> Modified: llvm/trunk/test/Transforms/InstCombine/2009-01-16-PointerAddrSpace.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/2009-01-16-PointerAddrSpace.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Transforms/InstCombine/2009-01-16-PointerAddrSpace.ll (original)
> +++ llvm/trunk/test/Transforms/InstCombine/2009-01-16-PointerAddrSpace.ll Thu Nov 14 19:34:59 2013
> @@ -5,7 +5,7 @@ target triple = "i386-apple-darwin9.6"
>
> define i32 @test(i32* %P) nounwind {
> entry:
> - %Q = bitcast i32* %P to i32 addrspace(1)*
> + %Q = addrspacecast i32* %P to i32 addrspace(1)*
> store i32 0, i32 addrspace(1)* %Q, align 4
> ret i32 0
> }
>
> Modified: llvm/trunk/test/Transforms/InstCombine/2012-07-30-addrsp-bitcast.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/2012-07-30-addrsp-bitcast.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Transforms/InstCombine/2012-07-30-addrsp-bitcast.ll (original)
> +++ llvm/trunk/test/Transforms/InstCombine/2012-07-30-addrsp-bitcast.ll Thu Nov 14 19:34:59 2013
> @@ -1,10 +1,10 @@
> ; RUN: opt < %s -instcombine -S | FileCheck %s
> -; CHECK: bitcast
> +; CHECK: addrspacecast
>
> @base = internal addrspace(3) unnamed_addr global [16 x i32] zeroinitializer, align 16
> declare void @foo(i32*)
>
> define void @test() nounwind {
> - call void @foo(i32* getelementptr (i32* bitcast ([16 x i32] addrspace(3)* @base to i32*), i64 2147483647)) nounwind
> + call void @foo(i32* getelementptr (i32* addrspacecast ([16 x i32] addrspace(3)* @base to i32*), i64 2147483647)) nounwind
> ret void
> }
>
> Modified: llvm/trunk/test/Transforms/InstCombine/cast_ptr.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/cast_ptr.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Transforms/InstCombine/cast_ptr.ll (original)
> +++ llvm/trunk/test/Transforms/InstCombine/cast_ptr.ll Thu Nov 14 19:34:59 2013
> @@ -100,7 +100,7 @@ define %unop* @test5(%op* %O) {
>
> define i8 @test6(i8 addrspace(1)* %source) {
> entry:
> - %arrayidx223 = bitcast i8 addrspace(1)* %source to i8*
> + %arrayidx223 = addrspacecast i8 addrspace(1)* %source to i8*
> %tmp4 = load i8* %arrayidx223
> ret i8 %tmp4
> ; CHECK-LABEL: @test6(
>
> 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=194760&r1=194759&r2=194760&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 Thu Nov 14 19:34:59 2013
> @@ -32,10 +32,10 @@ define i32 addrspace(2)* @test_constant_
> }
>
> ; Different address spaces that are the same size, but they are
> -; different so there should be a bitcast.
> +; different so nothing should happen
> define i32 addrspace(4)* @test_constant_fold_inttoptr_as_pointer_smaller_different_as() {
> ; CHECK-LABEL: @test_constant_fold_inttoptr_as_pointer_smaller_different_as(
> -; CHECK-NEXT: ret i32 addrspace(4)* bitcast (i32 addrspace(3)* @const_zero_i32_as3 to i32 addrspace(4)*)
> +; CHECK-NEXT: ret i32 addrspace(4)* inttoptr (i16 ptrtoint (i32 addrspace(3)* @const_zero_i32_as3 to i16) to i32 addrspace(4)*)
> %x = ptrtoint i32 addrspace(3)* @const_zero_i32_as3 to i16
> %y = inttoptr i16 %x to i32 addrspace(4)*
> ret i32 addrspace(4)* %y
>
> Modified: llvm/trunk/test/Transforms/InstCombine/gep-addrspace.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/gep-addrspace.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Transforms/InstCombine/gep-addrspace.ll (original)
> +++ llvm/trunk/test/Transforms/InstCombine/gep-addrspace.ll Thu Nov 14 19:34:59 2013
> @@ -9,7 +9,7 @@ target triple = "x86_64-pc-win32"
> define void @func(%myStruct addrspace(1)* nocapture %p) nounwind {
> ST:
> %A = getelementptr inbounds %myStruct addrspace(1)* %p, i64 0
> - %B = bitcast %myStruct addrspace(1)* %A to %myStruct*
> + %B = addrspacecast %myStruct addrspace(1)* %A to %myStruct*
> %C = getelementptr inbounds %myStruct* %B, i32 0, i32 1
> %D = getelementptr inbounds [3 x float]* %C, i32 0, i32 2
> %E = load float* %D, align 4
>
> Modified: llvm/trunk/test/Transforms/InstCombine/objsize-address-space.ll
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/InstCombine/objsize-address-space.ll?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/test/Transforms/InstCombine/objsize-address-space.ll (original)
> +++ llvm/trunk/test/Transforms/InstCombine/objsize-address-space.ll Thu Nov 14 19:34:59 2013
> @@ -49,7 +49,7 @@ define i32 @array_as2_size() {
> define i32 @pointer_array_as1() {
> ; CHECK-LABEL: @pointer_array_as1(
> ; CHECK-NEXT: ret i32 80
> - %bc = bitcast [10 x i32 addrspace(1)*]* @array_as1_pointers to i8 addrspace(1)*
> + %bc = addrspacecast [10 x i32 addrspace(1)*]* @array_as1_pointers to i8 addrspace(1)*
> %1 = call i32 @llvm.objectsize.i32.p1i8(i8 addrspace(1)* %bc, i1 false)
> ret i32 %1
> }
>
> Modified: llvm/trunk/unittests/IR/InstructionsTest.cpp
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/unittests/IR/InstructionsTest.cpp?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/unittests/IR/InstructionsTest.cpp (original)
> +++ llvm/trunk/unittests/IR/InstructionsTest.cpp Thu Nov 14 19:34:59 2013
> @@ -149,6 +149,8 @@ TEST(InstructionsTest, CastInst) {
> const Constant* c8 = Constant::getNullValue(V8x8Ty);
> const Constant* c64 = Constant::getNullValue(V8x64Ty);
>
> + const Constant *v2ptr32 = Constant::getNullValue(V2Int32PtrTy);
> +
> EXPECT_TRUE(CastInst::isCastable(V8x8Ty, X86MMXTy));
> EXPECT_TRUE(CastInst::isCastable(X86MMXTy, V8x8Ty));
> EXPECT_FALSE(CastInst::isCastable(Int64Ty, X86MMXTy));
> @@ -169,6 +171,10 @@ TEST(InstructionsTest, CastInst) {
> EXPECT_FALSE(CastInst::isBitCastable(V2Int32PtrTy, V2Int32PtrAS1Ty));
> EXPECT_FALSE(CastInst::isBitCastable(V2Int32PtrAS1Ty, V2Int32PtrTy));
> EXPECT_TRUE(CastInst::isBitCastable(V2Int32PtrAS1Ty, V2Int64PtrAS1Ty));
> + EXPECT_TRUE(CastInst::isCastable(V2Int32PtrAS1Ty, V2Int32PtrTy));
> + EXPECT_EQ(CastInst::AddrSpaceCast, CastInst::getCastOpcode(v2ptr32, true,
> + V2Int32PtrAS1Ty,
> + true));
>
> // Test mismatched number of elements for pointers
> EXPECT_FALSE(CastInst::isBitCastable(V2Int32PtrAS1Ty, V4Int64PtrAS1Ty));
> @@ -371,7 +377,6 @@ TEST(InstructionsTest, isEliminableCastP
> 0, Int32Ty, 0),
> 0U);
>
> -
> // Test that we don't eliminate bitcasts between different address spaces,
> // or if we don't have available pointer size information.
> DataLayout DL("e-p:32:32:32-p1:16:16:16-p2:64:64:64-i1:8:8-i8:8:8-i16:16:16"
> @@ -384,26 +389,18 @@ TEST(InstructionsTest, isEliminableCastP
> IntegerType *Int16SizePtr = DL.getIntPtrType(C, 1);
> IntegerType *Int64SizePtr = DL.getIntPtrType(C, 2);
>
> - // Fail since the ptr int types are not provided
> + // Cannot simplify inttoptr, addrspacecast
> EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::IntToPtr,
> - CastInst::BitCast,
> + CastInst::AddrSpaceCast,
> Int16Ty, Int64PtrTyAS1, Int64PtrTyAS2,
> - 0, 0, 0),
> - 0U);
> -
> - // Fail since the the bitcast is between different sized address spaces
> - EXPECT_EQ(CastInst::isEliminableCastPair(
> - CastInst::IntToPtr,
> - CastInst::BitCast,
> - Int16Ty, Int64PtrTyAS1, Int64PtrTyAS2,
> - 0, Int16SizePtr, Int64SizePtr),
> + 0, Int16SizePtr, Int64SizePtr),
> 0U);
>
> - // Fail since the the bitcast is between different sized address spaces
> - EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::IntToPtr,
> - CastInst::BitCast,
> - Int16Ty, Int64PtrTyAS1, Int64PtrTyAS2,
> - 0, Int16SizePtr, Int64SizePtr),
> + // Cannot simplify addrspacecast, ptrtoint
> + EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::AddrSpaceCast,
> + CastInst::PtrToInt,
> + Int64PtrTyAS1, Int64PtrTyAS2, Int16Ty,
> + Int64SizePtr, Int16SizePtr, 0),
> 0U);
>
> // Pass since the bitcast address spaces are the same
> @@ -413,28 +410,6 @@ TEST(InstructionsTest, isEliminableCastP
> 0, 0, 0),
> CastInst::IntToPtr);
>
> -
> - // Fail without known pointer sizes and different address spaces
> - EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::BitCast,
> - CastInst::PtrToInt,
> - Int64PtrTyAS1, Int64PtrTyAS2, Int16Ty,
> - 0, 0, 0),
> - 0U);
> -
> - // Pass since the address spaces are the same, even though the pointer sizes
> - // are unknown
> - EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::BitCast,
> - CastInst::PtrToInt,
> - Int64PtrTyAS1, Int64PtrTyAS1, Int32Ty,
> - 0, 0, 0),
> - Instruction::PtrToInt);
> -
> - // Fail since the bitcast is the wrong size
> - EXPECT_EQ(CastInst::isEliminableCastPair(CastInst::BitCast,
> - CastInst::PtrToInt,
> - Int64PtrTyAS1, Int64PtrTyAS2, Int64Ty,
> - Int16SizePtr, Int64SizePtr, 0),
> - 0U);
> }
>
> } // end anonymous namespace
>
> Modified: llvm/trunk/utils/emacs/llvm-mode.el
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/utils/emacs/llvm-mode.el?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/utils/emacs/llvm-mode.el (original)
> +++ llvm/trunk/utils/emacs/llvm-mode.el Thu Nov 14 19:34:59 2013
> @@ -43,7 +43,7 @@
> ;; Memory operators
> `(,(regexp-opt '("malloc" "alloca" "free" "load" "store" "getelementptr" "fence" "cmpxchg" "atomicrmw") 'words) . font-lock-keyword-face)
> ;; Casts
> - `(,(regexp-opt '("bitcast" "inttoptr" "ptrtoint" "trunc" "zext" "sext" "fptrunc" "fpext" "fptoui" "fptosi" "uitofp" "sitofp") 'words) . font-lock-keyword-face)
> + `(,(regexp-opt '("bitcast" "inttoptr" "ptrtoint" "trunc" "zext" "sext" "fptrunc" "fpext" "fptoui" "fptosi" "uitofp" "sitofp" "addrspacecast") 'words) . font-lock-keyword-face)
> ;; Vector ops
> `(,(regexp-opt '("extractelement" "insertelement" "shufflevector") 'words) . font-lock-keyword-face)
> ;; Aggregate ops
>
> Modified: llvm/trunk/utils/kate/llvm.xml
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/utils/kate/llvm.xml?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/utils/kate/llvm.xml (original)
> +++ llvm/trunk/utils/kate/llvm.xml Thu Nov 14 19:34:59 2013
> @@ -157,6 +157,7 @@
> <item> ptrtoint </item>
> <item> inttoptr </item>
> <item> bitcast </item>
> + <item> addrspacecast </item>
> <item> icmp </item>
> <item> fcmp </item>
> <item> phi </item>
>
> Modified: llvm/trunk/utils/vim/llvm.vim
> URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/utils/vim/llvm.vim?rev=194760&r1=194759&r2=194760&view=diff
> ==============================================================================
> --- llvm/trunk/utils/vim/llvm.vim (original)
> +++ llvm/trunk/utils/vim/llvm.vim Thu Nov 14 19:34:59 2013
> @@ -22,17 +22,18 @@ syn match llvmType /\<i\d\+\>/
> " Instructions.
> " The true and false tokens can be used for comparison opcodes, but it's
> " much more common for these tokens to be used for boolean constants.
> -syn keyword llvmStatement add alloca and arcp ashr atomicrmw bitcast br call
> -syn keyword llvmStatement cmpxchg eq exact extractelement extractvalue fadd fast
> -syn keyword llvmStatement fcmp fdiv fence fmul fpext fptosi fptoui fptrunc free
> -syn keyword llvmStatement frem fsub getelementptr icmp inbounds indirectbr
> -syn keyword llvmStatement insertelement insertvalue inttoptr invoke landingpad
> -syn keyword llvmStatement load lshr malloc max min mul nand ne ninf nnan nsw nsz
> -syn keyword llvmStatement nuw oeq oge ogt ole olt one or ord phi ptrtoint resume
> -syn keyword llvmStatement ret sdiv select sext sge sgt shl shufflevector sitofp
> -syn keyword llvmStatement sle slt srem store sub switch trunc udiv ueq uge ugt
> -syn keyword llvmStatement uitofp ule ult umax umin une uno unreachable unwind
> -syn keyword llvmStatement urem va_arg xchg xor zext
> +syn keyword llvmStatement add addrspacecast alloca and arcp ashr atomicrmw
> +syn keyword llvmStatement bitcast br call cmpxchg eq exact extractelement
> +syn keyword llvmStatement extractvalue fadd fast fcmp fdiv fence fmul fpext
> +syn keyword llvmStatement fptosi fptoui fptrunc free frem fsub getelementptr
> +syn keyword llvmStatement icmp inbounds indirectbr insertelement insertvalue
> +syn keyword llvmStatement inttoptr invoke landingpad load lshr malloc max min
> +syn keyword llvmStatement mul nand ne ninf nnan nsw nsz nuw oeq oge ogt ole
> +syn keyword llvmStatement olt one or ord phi ptrtoint resume ret sdiv select
> +syn keyword llvmStatement sext sge sgt shl shufflevector sitofp sle slt srem
> +syn keyword llvmStatement store sub switch trunc udiv ueq uge ugt uitofp ule ult
> +syn keyword llvmStatement umax umin une uno unreachable unwind urem va_arg
> +syn keyword llvmStatement xchg xor zext
>
> " Keywords.
> syn keyword llvmKeyword acq_rel acquire sanitize_address addrspace alias align
>
More information about the llvm-commits
mailing list