Merge branch 'amd-master' into amd-common

cmake/modules/HandleLLVMOptions.cmake

@@ -475,7 +475,7 @@ elseif( LLVM_COMPILER_IS_GCC_COMPATIBLE )
   endif()
   if (LLVM_ENABLE_MODULES)
     set(OLD_CMAKE_REQUIRED_FLAGS ${CMAKE_REQUIRED_FLAGS})
-    set(module_flags "-fmodules -fmodules-cache-path=module.cache")
+    set(module_flags "-fmodules -fmodules-cache-path=${PROJECT_BINARY_DIR}/module.cache")
     if (${CMAKE_SYSTEM_NAME} MATCHES "Darwin")
       # On Darwin -fmodules does not imply -fcxx-modules.
       set(module_flags "${module_flags} -fcxx-modules")

docs/HowToAddABuilder.rst

@@ -30,7 +30,7 @@ Here are the steps you can follow to do so:
 
 #. Install buildslave (currently we are using buildbot version 0.8.5).
    Depending on the platform, buildslave could be available to download and
-   install with your packet manager, or you can download it directly from
+   install with your package manager, or you can download it directly from
    `<http://trac.buildbot.net>`_ and install it manually.
 
 #. Create a designated user account, your buildslave will be running under,
@@ -43,7 +43,7 @@ Here are the steps you can follow to do so:
 #. Create a buildslave in context of that buildslave account.  Point it to
    the **lab.llvm.org** port **9990** (see `Buildbot documentation,
    Creating a slave
-   <http://buildbot.net/buildbot/docs/current/full.html#creating-a-slave>`_
+   <http://docs.buildbot.net/current/tutorial/firstrun.html#creating-a-slave>`_
    for more details) by running the following command:
 
     .. code-block:: bash

include/llvm/Analysis/ProfileSummaryInfo.h

@@ -69,7 +69,12 @@ class ProfileSummaryInfoWrapperPass : public ImmutablePass {
   static char ID;
   ProfileSummaryInfoWrapperPass();
 
-  ProfileSummaryInfo *getPSI(Module &M);
+  ProfileSummaryInfo *getPSI() {
+    return &*PSI;
+  }
+
+  bool doInitialization(Module &M) override;
+  bool doFinalization(Module &M) override;
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.setPreservesAll();
   }

include/llvm/Analysis/ScalarEvolution.h

@@ -551,19 +551,36 @@ class ScalarEvolution {
     const SCEV *ExactNotTaken;
     const SCEV *MaxNotTaken;
 
-    /// A predicate union guard for this ExitLimit. The result is only
-    /// valid if this predicate evaluates to 'true' at run-time.
-    SCEVUnionPredicate Predicate;
+    /// A set of predicate guards for this ExitLimit. The result is only valid
+    /// if all of the predicates in \c Predicates evaluate to 'true' at
+    /// run-time.
+    SmallPtrSet<const SCEVPredicate *, 4> Predicates;
+
+    void addPredicate(const SCEVPredicate *P) {
+      assert(!isa<SCEVUnionPredicate>(P) && "Only add leaf predicates here!");
+      Predicates.insert(P);
+    }
 
     /*implicit*/ ExitLimit(const SCEV *E) : ExactNotTaken(E), MaxNotTaken(E) {}
 
-    ExitLimit(const SCEV *E, const SCEV *M, SCEVUnionPredicate &P)
-        : ExactNotTaken(E), MaxNotTaken(M), Predicate(P) {
+    ExitLimit(
+        const SCEV *E, const SCEV *M,
+        ArrayRef<const SmallPtrSetImpl<const SCEVPredicate *> *> PredSetList)
+        : ExactNotTaken(E), MaxNotTaken(M) {
       assert((isa<SCEVCouldNotCompute>(ExactNotTaken) ||
               !isa<SCEVCouldNotCompute>(MaxNotTaken)) &&
              "Exact is not allowed to be less precise than Max");
+      for (auto *PredSet : PredSetList)
+        for (auto *P : *PredSet)
+          addPredicate(P);
     }
 
+    ExitLimit(const SCEV *E, const SCEV *M,
+              const SmallPtrSetImpl<const SCEVPredicate *> &PredSet)
+        : ExitLimit(E, M, {&PredSet}) {}
+
+    ExitLimit(const SCEV *E, const SCEV *M) : ExitLimit(E, M, None) {}
+
     /// Test whether this ExitLimit contains any computed information, or
     /// whether it's all SCEVCouldNotCompute values.
     bool hasAnyInfo() const {
@@ -1581,9 +1598,9 @@ class ScalarEvolution {
                                     SCEVUnionPredicate &A);
   /// Tries to convert the \p S expression to an AddRec expression,
   /// adding additional predicates to \p Preds as required.
-  const SCEVAddRecExpr *
-  convertSCEVToAddRecWithPredicates(const SCEV *S, const Loop *L,
-                                    SCEVUnionPredicate &Preds);
+  const SCEVAddRecExpr *convertSCEVToAddRecWithPredicates(
+      const SCEV *S, const Loop *L,
+      SmallPtrSetImpl<const SCEVPredicate *> &Preds);
 
 private:
   /// Compute the backedge taken count knowing the interval difference, the

include/llvm/CodeGen/GlobalISel/RegisterBankInfo.h

@@ -143,10 +143,23 @@ class RegisterBankInfo {
     /// Number of partial mapping to break down this value.
     unsigned NumBreakDowns;
 
+    /// The default constructor creates an invalid (isValid() == false)
+    /// instance.
+    ValueMapping() : ValueMapping(nullptr, 0) {}
+
+    /// Initialize a ValueMapping with the given parameter.
+    /// \p BreakDown needs to have a life time at least as long
+    /// as this instance.
+    ValueMapping(const PartialMapping *BreakDown, unsigned NumBreakDowns)
+        : BreakDown(BreakDown), NumBreakDowns(NumBreakDowns) {}
+
     /// Iterators through the PartialMappings.
     const PartialMapping *begin() const { return BreakDown; }
     const PartialMapping *end() const { return BreakDown + NumBreakDowns; }
 
+    /// Check if this ValueMapping is valid.
+    bool isValid() const { return BreakDown && NumBreakDowns; }
+
     /// Verify that this mapping makes sense for a value of
     /// \p MeaningFulBitWidth.
     /// \note This method does not check anything when assertions are disabled.
@@ -171,12 +184,11 @@ class RegisterBankInfo {
     /// Cost of this mapping.
     unsigned Cost;
     /// Mapping of all the operands.
-    /// Note: Use a SmallVector to avoid heap allocation in most cases.
-    SmallVector<const ValueMapping *, 8> OperandsMapping;
+    const ValueMapping *OperandsMapping;
     /// Number of operands.
     unsigned NumOperands;
 
-    const ValueMapping *&getOperandMapping(unsigned i) {
+    const ValueMapping &getOperandMapping(unsigned i) {
       assert(i < getNumOperands() && "Out of bound operand");
       return OperandsMapping[i];
     }
@@ -190,11 +202,13 @@ class RegisterBankInfo {
     /// at the index i.
     ///
     /// \pre ID != InvalidMappingID
-    InstructionMapping(unsigned ID, unsigned Cost, unsigned NumOperands)
-        : ID(ID), Cost(Cost), NumOperands(NumOperands) {
+    InstructionMapping(unsigned ID, unsigned Cost,
+                       const ValueMapping *OperandsMapping,
+                       unsigned NumOperands)
+        : ID(ID), Cost(Cost), OperandsMapping(OperandsMapping),
+          NumOperands(NumOperands) {
       assert(getID() != InvalidMappingID &&
              "Use the default constructor for invalid mapping");
-      OperandsMapping.resize(getNumOperands(), nullptr);
     }
 
     /// Default constructor.
@@ -214,26 +228,23 @@ class RegisterBankInfo {
     /// \pre The mapping for the ith operand has been set.
     /// \pre The ith operand is a register.
     const ValueMapping &getOperandMapping(unsigned i) const {
-      const ValueMapping *&ValMapping =
+      const ValueMapping &ValMapping =
           const_cast<InstructionMapping *>(this)->getOperandMapping(i);
-      assert(ValMapping && "Trying to get the mapping for a non-reg operand?");
-      return *ValMapping;
-    }
-
-    /// Check if the value mapping of the ith operand has been set.
-    bool isOperandMappingSet(unsigned i) const {
-      return const_cast<InstructionMapping *>(this)->getOperandMapping(i) !=
-             nullptr;
+      return ValMapping;
     }
 
-    /// Get the value mapping of the ith operand.
-    void setOperandMapping(unsigned i, const ValueMapping &ValMapping) {
-      getOperandMapping(i) = &ValMapping;
+    /// Set the mapping for all the operands.
+    /// In other words, OpdsMapping should hold at least getNumOperands
+    /// ValueMapping.
+    void setOperandsMapping(const ValueMapping *OpdsMapping) {
+      OperandsMapping = OpdsMapping;
     }
 
     /// Check whether this object is valid.
     /// This is a lightweight check for obvious wrong instance.
-    bool isValid() const { return getID() != InvalidMappingID; }
+    bool isValid() const {
+      return getID() != InvalidMappingID && OperandsMapping;
+    }
 
     /// Verifiy that this mapping makes sense for \p MI.
     /// \pre \p MI must be connected to a MachineFunction.
@@ -367,6 +378,10 @@ class RegisterBankInfo {
   /// This shouldn't be needed when everything gets TableGen'ed.
   mutable DenseMap<unsigned, const ValueMapping *> MapOfValueMappings;
 
+  /// Keep dynamically allocated array of ValueMapping in a separate map.
+  /// This shouldn't be needed when everything gets TableGen'ed.
+  mutable DenseMap<unsigned, ValueMapping *> MapOfOperandsMappings;
+
   /// Create a RegisterBankInfo that can accomodate up to \p NumRegBanks
   /// RegisterBank instances.
   ///
@@ -453,6 +468,39 @@ class RegisterBankInfo {
                                       unsigned NumBreakDowns) const;
   /// @}
 
+  /// Methods to get a uniquely generated array of ValueMapping.
+  /// @{
+
+  /// Get the uniquely generated array of ValueMapping for the
+  /// elements of between \p Begin and \p End.
+  ///
+  /// Elements that are nullptr will be replaced by
+  /// invalid ValueMapping (ValueMapping::isValid == false).
+  ///
+  /// \pre The pointers on ValueMapping between \p Begin and \p End
+  /// must uniquely identify a ValueMapping. Otherwise, there is no
+  /// guarantee that the return instance will be unique, i.e., another
+  /// OperandsMapping could have the same content.
+  template <typename Iterator>
+  const ValueMapping *getOperandsMapping(Iterator Begin, Iterator End) const;
+
+  /// Get the uniquely generated array of ValueMapping for the
+  /// elements of \p OpdsMapping.
+  ///
+  /// Elements of \p OpdsMapping that are nullptr will be replaced by
+  /// invalid ValueMapping (ValueMapping::isValid == false).
+  const ValueMapping *getOperandsMapping(
+      const SmallVectorImpl<const ValueMapping *> &OpdsMapping) const;
+
+  /// Get the uniquely generated array of ValueMapping for the
+  /// given arguments.
+  ///
+  /// Arguments that are nullptr will be replaced by invalid
+  /// ValueMapping (ValueMapping::isValid == false).
+  const ValueMapping *getOperandsMapping(
+      std::initializer_list<const ValueMapping *> OpdsMapping) const;
+  /// @}
+
   /// Get the register bank for the \p OpIdx-th operand of \p MI form
   /// the encoding constraints, if any.
   ///

include/llvm/CodeGen/ScheduleDAGInstrs.h

@@ -345,7 +345,6 @@ namespace llvm {
     SUnits.emplace_back(MI, (unsigned)SUnits.size());
     assert((Addr == nullptr || Addr == &SUnits[0]) &&
            "SUnits std::vector reallocated on the fly!");
-    SUnits.back().OrigNode = &SUnits.back();
     return &SUnits.back();
   }
 

include/llvm/DebugInfo/DIContext.h

@@ -42,6 +42,10 @@ struct DILineInfo {
   bool operator!=(const DILineInfo &RHS) const {
     return !(*this == RHS);
   }
+  bool operator<(const DILineInfo &RHS) const {
+    return std::tie(FileName, FunctionName, Line, Column) <
+           std::tie(RHS.FileName, RHS.FunctionName, RHS.Line, RHS.Column);
+  }
 };
 
 typedef SmallVector<std::pair<uint64_t, DILineInfo>, 16> DILineInfoTable;

include/llvm/IR/Attributes.h

@@ -68,7 +68,7 @@ class Attribute {
     // IR-Level Attributes
     None,                  ///< No attributes have been set
     #define GET_ATTR_ENUM
-    #include "llvm/IR/Attributes.inc"
+    #include "llvm/IR/Attributes.gen"
     EndAttrKinds           ///< Sentinal value useful for loops
   };
 

include/llvm/IR/CMakeLists.txt

@@ -1,5 +1,5 @@
 set(LLVM_TARGET_DEFINITIONS Attributes.td)
-tablegen(LLVM Attributes.inc -gen-attrs)
+tablegen(LLVM Attributes.gen -gen-attrs)
 
 set(LLVM_TARGET_DEFINITIONS Intrinsics.td)
 tablegen(LLVM Intrinsics.gen -gen-intrinsic)

include/llvm/IR/IntrinsicsNVVM.td

@@ -729,6 +729,39 @@ let TargetPrefix = "nvvm" in {
           [LLVMAnyPointerType<llvm_i32_ty>, llvm_i32_ty],
                                       [IntrArgMemOnly, NoCapture<0>]>;
 
+  class SCOPED_ATOMIC2_impl<LLVMType elty>
+        : Intrinsic<[elty],
+          [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>],
+          [IntrArgMemOnly, NoCapture<0>]>;
+  class SCOPED_ATOMIC3_impl<LLVMType elty>
+        : Intrinsic<[elty],
+          [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>,
+           LLVMMatchType<0>],
+          [IntrArgMemOnly, NoCapture<0>]>;
+
+  multiclass PTXAtomicWithScope2<LLVMType elty> {
+    def _cta : SCOPED_ATOMIC2_impl<elty>;
+    def _sys : SCOPED_ATOMIC2_impl<elty>;
+  }
+  multiclass PTXAtomicWithScope3<LLVMType elty> {
+    def _cta : SCOPED_ATOMIC3_impl<elty>;
+    def _sys : SCOPED_ATOMIC3_impl<elty>;
+  }
+  multiclass PTXAtomicWithScope2_fi {
+    defm _f: PTXAtomicWithScope2<llvm_anyfloat_ty>;
+    defm _i: PTXAtomicWithScope2<llvm_anyint_ty>;
+  }
+  defm int_nvvm_atomic_add_gen   : PTXAtomicWithScope2_fi;
+  defm int_nvvm_atomic_inc_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_dec_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_exch_gen_i: PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_xor_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_max_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_min_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_or_gen_i  : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_and_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
+  defm int_nvvm_atomic_cas_gen_i : PTXAtomicWithScope3<llvm_anyint_ty>;
+
 // Bar.Sync
 
   // The builtin for "bar.sync 0" is called __syncthreads.  Unlike most of the

include/llvm/IR/Mangler.h

@@ -29,12 +29,7 @@ class Mangler {
   /// This keeps track of the number we give to anonymous ones.
   mutable DenseMap<const GlobalValue*, unsigned> AnonGlobalIDs;
 
-  /// This simple counter is used to unique value names.
-  mutable unsigned NextAnonGlobalID;
-
 public:
-  Mangler() : NextAnonGlobalID(1) {}
-
   /// Print the appropriate prefix and the specified global variable's name.
   /// If the global variable doesn't have a name, this fills in a unique name
   /// for the global.

include/llvm/LTO/LTO.h

@@ -175,6 +175,8 @@ class InputFile {
       // FIXME: Expose a thread-local flag for module asm symbols.
       return GV && GV->isThreadLocal();
     }
+
+    //FIXME: We shouldn't expose this information.
     Expected<const Comdat *> getComdat() const {
       if (!GV)
         return nullptr;
@@ -191,6 +193,7 @@ class InputFile {
         return GO->getComdat();
       return nullptr;
     }
+
     uint64_t getCommonSize() const {
       assert(Flags & object::BasicSymbolRef::SF_Common);
       if (!GV)
@@ -249,6 +252,11 @@ class InputFile {
   MemoryBufferRef getMemoryBufferRef() const {
     return Obj->getMemoryBufferRef();
   }
+
+  // FIXME: We should fix lld and not expose this information.
+  StringMap<Comdat> &getComdatSymbolTable() {
+    return Obj->getModule().getComdatSymbolTable();
+  }
 };
 
 /// This class wraps an output stream for a native object. Most clients should