From bd722653dbfb7f2468c4facec82e63b3f82a10d6 Mon Sep 17 00:00:00 2001
From: Wolfgang Puffitsch <hausen@gmx.at>
Date: Fri, 15 Apr 2016 08:59:43 +0200
Subject: [PATCH] Make things work again after upgrade to LLVM 3.8.

---
 include/clang-c/Index.h                       |    6 +-
 include/clang/AST/RecursiveASTVisitor.h       |    3 +
 include/clang/AST/StmtPlatin.h                |    4 +-
 include/clang/Basic/Attr.td                   |   28 +-
 include/clang/Basic/Builtins.def              |    3 +
 include/clang/Basic/Builtins.h                |    2 +-
 include/clang/Basic/DiagnosticParseKinds.td   |    8 +
 include/clang/Basic/DiagnosticSemaKinds.td    |    2 +
 include/clang/Basic/StmtNodes.td              |    3 +
 include/clang/Basic/TokenKinds.def            |    9 +
 include/clang/Driver/Driver.h                 |    2 +-
 include/clang/Driver/Options.td               |   78 +-
 include/clang/Driver/ToolChain.h              |   10 +-
 include/clang/Frontend/ASTConsumers.h         |    2 +-
 include/clang/Frontend/FrontendActions.h      |    4 +-
 include/clang/Frontend/FrontendOptions.h      |    3 +
 include/clang/Parse/Parser.h                  |   13 +
 include/clang/Sema/Sema.h                     |    6 +
 lib/AST/Stmt.cpp                              |   32 +
 lib/AST/StmtPrinter.cpp                       |    9 +
 lib/AST/StmtProfile.cpp                       |    4 +
 lib/Basic/Targets.cpp                         |  151 ++
 lib/CodeGen/BackendUtil.cpp                   |    6 -
 lib/CodeGen/CGBuiltin.cpp                     |    8 +-
 lib/CodeGen/CGStmt.cpp                        |  108 +-
 lib/CodeGen/CodeGenFunction.h                 |    9 +-
 lib/CodeGen/TargetInfo.cpp                    |   27 +
 lib/Driver/CMakeLists.txt                     |    1 +
 lib/Driver/Driver.cpp                         |   11 +-
 lib/Driver/ToolChains.cpp                     |  127 ++
 lib/Driver/ToolChains.h                       |   37 +
 lib/Driver/Tools.cpp                          | 1352 +++++++++++++++++
 lib/Driver/Tools.h                            |  161 ++
 lib/Frontend/CompilerInvocation.cpp           |    1 +
 lib/Frontend/FrontendActions.cpp              |    7 +-
 lib/Parse/ParsePragma.cpp                     |  140 ++
 lib/Parse/ParseStmt.cpp                       |   36 +
 lib/Parse/Parser.cpp                          |    6 -
 lib/Sema/SemaStmtAttr.cpp                     |   35 +
 lib/Sema/TreeTransform.h                      |   16 +
 lib/Serialization/ASTReaderStmt.cpp           |    8 +
 lib/Serialization/ASTWriterStmt.cpp           |    5 +
 lib/Serialization/CMakeLists.txt              |    1 +
 lib/Serialization/FlowfactExporter.cpp        |   14 +-
 lib/StaticAnalyzer/Core/ExprEngine.cpp        |    1 +
 test/CodeGenCUDA/address-space-conversion.cu  |   48 -
 .../CodeGenOpenCL/address-space-conversion.cl |   14 -
 tools/libclang/CIndex.cpp                     |    2 +
 tools/libclang/CXCursor.cpp                   |    3 +
 49 files changed, 2415 insertions(+), 151 deletions(-)
 delete mode 100644 test/CodeGenCUDA/address-space-conversion.cu
 delete mode 100644 test/CodeGenOpenCL/address-space-conversion.cl

diff --git a/include/clang-c/Index.h b/include/clang-c/Index.h
index 09e2160826..ed4f4419d5 100644
--- a/include/clang-c/Index.h
+++ b/include/clang-c/Index.h
@@ -2274,7 +2274,11 @@ enum CXCursorKind {
    */
   CXCursor_OMPDistributeDirective        = 260,
 
-  CXCursor_LastStmt                      = CXCursor_OMPDistributeDirective,
+   /** \brief Flowfact sepcification.
+   */
+  CXCursor_Flowfact                      = 261,
+
+  CXCursor_LastStmt                      = CXCursor_Flowfact,
 
   /**
    * \brief Cursor that represents the translation unit itself.
diff --git a/include/clang/AST/RecursiveASTVisitor.h b/include/clang/AST/RecursiveASTVisitor.h
index 4c2c43d610..41877c5ccc 100644
--- a/include/clang/AST/RecursiveASTVisitor.h
+++ b/include/clang/AST/RecursiveASTVisitor.h
@@ -2341,6 +2341,9 @@ DEF_TRAVERSE_STMT(ObjCBoxedExpr, {})
 DEF_TRAVERSE_STMT(ObjCArrayLiteral, {})
 DEF_TRAVERSE_STMT(ObjCDictionaryLiteral, {})
 
+// Traverse Flowfact
+DEF_TRAVERSE_STMT(Flowfact, {})
+
 // Traverse OpenCL: AsType, Convert.
 DEF_TRAVERSE_STMT(AsTypeExpr, {})
 
diff --git a/include/clang/AST/StmtPlatin.h b/include/clang/AST/StmtPlatin.h
index 419556b201..f0cfda9f50 100644
--- a/include/clang/AST/StmtPlatin.h
+++ b/include/clang/AST/StmtPlatin.h
@@ -26,8 +26,8 @@ public:
   /// \brief Returns ending location of the flowfact.
   SourceLocation getLocEnd() const { return Range.getEnd(); }
 
-  StmtRange children() {
-    return StmtRange();
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
   }
 
   unsigned getNumLhsTerms() const { return Multipliers.size(); }
diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index 30381aa3d9..517fd1967f 100644
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -258,6 +258,7 @@ def TargetWindows : TargetArch<["x86", "x86_64", "arm", "thumb"]> {
 def TargetMicrosoftCXXABI : TargetArch<["x86", "x86_64", "arm", "thumb"]> {
   let CXXABIs = ["Microsoft"];
 }
+def TargetSinglePath : TargetArch<["patmos"]>;
 
 class Attr {
   // The various ways in which an attribute can be spelled in source
@@ -1320,9 +1321,10 @@ def Sentinel : InheritableAttr {
 }
 
 // For targets that support single-path code generation
-def SinglePath : InheritableAttr, TargetSpecificAttr {
+def SinglePath : InheritableAttr, TargetSpecificAttr<TargetSinglePath> {
   let Spellings = [GNU<"singlepath">, CXX11<"gnu", "singlepath">];
-  let Subjects = [Function];
+  let Subjects = SubjectList<[Function]>;
+  let Documentation = [Undocumented];
 }
 
 def StdCall : InheritableAttr {
@@ -2157,6 +2159,28 @@ def LoopHint : Attr {
   let Documentation = [LoopHintDocs, UnrollHintDocs];
 }
 
+def Loopbound : Attr {
+  // #pragma loopbound min VALUE max VALUE
+  let Spellings = [Pragma<"", "loopbound">];
+  let Args = [DefaultIntArgument<"Min", -1>,
+              DefaultIntArgument<"Max", -1>];
+  let Documentation = [Undocumented];
+
+  let AdditionalMembers = [{
+    void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const {
+      OS << "min " << getValueString(min, Policy)
+         << "max " << getValueString(max, Policy) << "\n";
+     }
+
+    std::string getValueString(int val, const PrintingPolicy &Policy) const {
+      std::string ValueName;
+      llvm::raw_string_ostream OS(ValueName);
+      OS << val << " ";
+      return OS.str();
+    }
+  }];
+}
+
 def CapturedRecord : InheritableAttr {
   // This attribute has no spellings as it is only ever created implicitly.
   let Spellings = [];
diff --git a/include/clang/Basic/Builtins.def b/include/clang/Basic/Builtins.def
index 4f474ebe42..7eee558b2e 100644
--- a/include/clang/Basic/Builtins.def
+++ b/include/clang/Basic/Builtins.def
@@ -718,6 +718,9 @@ LANGBUILTIN(__va_start,       "vc**.", "nt", ALL_MS_LANGUAGES)
 // Microsoft library builtins.
 LIBBUILTIN(_setjmpex, "iJ", "fj",   "setjmpex.h", ALL_MS_LANGUAGES)
 
+// Builtins for LLVM intrinsics
+BUILTIN(__llvm_pcmarker, "vi", "n")
+
 // C99 library functions
 // C99 stdlib.h
 LIBBUILTIN(abort, "v",            "fr",    "stdlib.h", ALL_LANGUAGES)
diff --git a/include/clang/Basic/Builtins.h b/include/clang/Basic/Builtins.h
index 24513cade1..75d028eeb7 100644
--- a/include/clang/Basic/Builtins.h
+++ b/include/clang/Basic/Builtins.h
@@ -143,7 +143,7 @@ public:
   /// \brief Return true if this is a builtin that should translate
   /// to a LLVM intrinsic of the same name
   bool isLLVMIntrinsicFunction(unsigned ID) const {
-    return strncmp("__llvm_", GetRecord(ID).Name, 7) == 0;
+    return strncmp("__llvm_", getRecord(ID).Name, 7) == 0;
   }
 
   /// \brief Determines whether this builtin has custom typechecking.
diff --git a/include/clang/Basic/DiagnosticParseKinds.td b/include/clang/Basic/DiagnosticParseKinds.td
index f8dee2f98c..fbed8c73d6 100644
--- a/include/clang/Basic/DiagnosticParseKinds.td
+++ b/include/clang/Basic/DiagnosticParseKinds.td
@@ -982,4 +982,12 @@ def err_for_co_await_not_range_for : Error<
   "'co_await' modifier can only be applied to range-based for loop">;
 }
 
+// - #pragma loopbound
+def err_pragma_loopbound_malformed : Error<
+  "pragma loopbound is malformed; expecting "
+  "'#pragma loopbound min NUM max NUM'">;
+
+def err_pragma_platin_malformed : Error<
+  "platin pragma needs to be flowfact: (c1 @x + c2 @y <= c3)">;
+
 } // end of Parser diagnostics
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 6ba482c78e..97817af35f 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -660,6 +660,8 @@ def err_pragma_loop_compatibility : Error<
   "%select{incompatible|duplicate}0 directives '%1' and '%2'">;
 def err_pragma_loop_precedes_nonloop : Error<
   "expected a for, while, or do-while loop to follow '%0'">;
+def err_pragma_loopbound_invalid_values : Error<
+  "invalid values; expected min >= 0 && max >= 0 && min <= max">;
 
 /// Objective-C parser diagnostics
 def err_duplicate_class_def : Error<
diff --git a/include/clang/Basic/StmtNodes.td b/include/clang/Basic/StmtNodes.td
index 36519ea29c..1972efc3ae 100644
--- a/include/clang/Basic/StmtNodes.td
+++ b/include/clang/Basic/StmtNodes.td
@@ -191,6 +191,9 @@ def MSDependentExistsStmt : Stmt;
 // OpenCL Extensions.
 def AsTypeExpr : DStmt<Expr>;
 
+// Flow facts
+def Flowfact : Stmt;
+
 // OpenMP Directives.
 def OMPExecutableDirective : Stmt<1>;
 def OMPLoopDirective : DStmt<OMPExecutableDirective, 1>;
diff --git a/include/clang/Basic/TokenKinds.def b/include/clang/Basic/TokenKinds.def
index 026945141d..02c4305303 100644
--- a/include/clang/Basic/TokenKinds.def
+++ b/include/clang/Basic/TokenKinds.def
@@ -765,6 +765,15 @@ ANNOTATION(pragma_openmp_end)
 // handles #pragma loop ... directives.
 ANNOTATION(pragma_loop_hint)
 
+// Annotations for #pragma loopbound
+// The lexer produces these so that they only take effect when the parser
+// handles #pragma loopbound ... directives.
+ANNOTATION(pragma_loopbound)
+
+// Annotations for platin #pragmas
+ANNOTATION(pragma_platinff)
+ANNOTATION(pragma_platinff_end)
+
 // Annotations for module import translated from #include etc.
 ANNOTATION(module_include)
 ANNOTATION(module_begin)
diff --git a/include/clang/Driver/Driver.h b/include/clang/Driver/Driver.h
index a229779e1a..1f09563a12 100644
--- a/include/clang/Driver/Driver.h
+++ b/include/clang/Driver/Driver.h
@@ -426,7 +426,7 @@ public:
 private:
   /// Parse the \p Args list for LTO options and record the type of LTO
   /// compilation based on which -f(no-)?lto(=.*)? option occurs last.
-  void setLTOMode(const llvm::opt::ArgList &Args);
+  void setLTOMode(const ToolChain &TC, const llvm::opt::ArgList &Args);
 
   /// \brief Retrieves a ToolChain for a particular \p Target triple.
   ///
diff --git a/include/clang/Driver/Options.td b/include/clang/Driver/Options.td
index 5da9602478..b223c552ab 100644
--- a/include/clang/Driver/Options.td
+++ b/include/clang/Driver/Options.td
@@ -95,6 +95,11 @@ def m_ppc_Features_Group : OptionGroup<"<ppc features group>">,
                            Group<m_Group>;
 def m_wasm_Features_Group : OptionGroup<"<wasm features group>">,
                             Group<m_Group>;
+def m_patmos_Features_Group : OptionGroup<"<m Patmos features group>">,
+                              Group<m_Group>;
+def m_patmos_llc_Group    : OptionGroup<"<m Patmos llc group>">;
+
+def m_PML_Group           : OptionGroup<"<m PML group>">;
 
 def m_libc_Group          : OptionGroup<"<m libc group>">, Group<m_Group>;
 def u_Group               : OptionGroup<"<u group>">;
@@ -340,6 +345,8 @@ def z : Separate<["-"], "z">, Flags<[LinkerInput, RenderAsInput]>,
   HelpText<"Pass -z <arg> to the linker">, MetaVarName<"<arg>">;
 def Xlinker : Separate<["-"], "Xlinker">, Flags<[LinkerInput, RenderAsInput]>,
   HelpText<"Pass <arg> to the linker">, MetaVarName<"<arg>">;
+def Xpreprocessor : Separate<["-"], "Xpreprocessor">,
+  HelpText<"Pass <arg> to the preprocessor">, MetaVarName<"<arg>">;
 def Xopt : Separate<["-"], "Xopt">, Flags<[RenderAsInput]>,
   HelpText<"Pass <arg> to the bitcode optimizer (Patmos only)">, 
   MetaVarName<"<arg>">;
@@ -349,8 +356,6 @@ def Xllc : Separate<["-"], "Xllc">, Flags<[RenderAsInput]>,
 def Xgold : Separate<["-"], "Xgold">, Flags<[LinkerInput, RenderAsInput]>,
   HelpText<"Pass <arg> to the gold ELF linker (Patmos only)">, 
   MetaVarName<"<arg>">;
-def Xpreprocessor : Separate<["-"], "Xpreprocessor">,
-  HelpText<"Pass <arg> to the preprocessor">, MetaVarName<"<arg>">;
 def X_Flag : Flag<["-"], "X">;
 def X_Joined : Joined<["-"], "X">;
 def Z_Flag : Flag<["-"], "Z">;
@@ -992,17 +997,17 @@ def fpascal_strings : Flag<["-"], "fpascal-strings">, Group<f_Group>, Flags<[CC1
 
 def fpatmos_emit_llvm : Flag<["-"], "fpatmos-emit-llvm">,
   HelpText<"Generate bitcode instead of a native binary">;
-def fpatmos_emit_reloc : Flag<["-"], "fpatmos-emit-reloc">, 
+def fpatmos_emit_reloc : Flag<["-"], "fpatmos-emit-reloc">,
   HelpText<"Generate a relocatable object instead of an executable or bitcode">;
 def fpatmos_emit_asm : Flag<["-"], "fpatmos-emit-asm">,
   HelpText<"Generate native assembler code">;
-def fpatmos_link_object: Flag<["-"], "fpatmos-link-object">, 
+def fpatmos_link_object: Flag<["-"], "fpatmos-link-object">,
   HelpText<"Link as (relocatable) object, i.e., do not link in libraries or startup code. Emits bitcode by default.">;
-def fpatmos_disable_internalize: Flag<["-"], "fpatmos-disable-internalize">, 
+def fpatmos_disable_internalize: Flag<["-"], "fpatmos-disable-internalize">,
   HelpText<"Disable internalize pass after bitcode linking">;
 def fpatmos_no_std_link_opts: Flag<["-"], "fpatmos-no-std-link-opts">,
   HelpText<"Disable default opt passes after bitcode linking">;
-def fpatmos_skip_opt: Flag<["-"], "fpatmos-skip-opt">, 
+def fpatmos_skip_opt: Flag<["-"], "fpatmos-skip-opt">,
   HelpText<"Skip opt phase after bitcode linker">;
 
 def fpcc_struct_return : Flag<["-"], "fpcc-struct-return">, Group<f_Group>, Flags<[CC1Option]>,
@@ -1651,61 +1656,61 @@ def multiply__defined__unused : Separate<["-"], "multiply_defined_unused">;
 def multiply__defined : Separate<["-"], "multiply_defined">;
 def mwarn_nonportable_cfstrings : Flag<["-"], "mwarn-nonportable-cfstrings">, Group<m_Group>;
 
-def mpatmos_stack_base : Joined<["-"], "mpatmos-stack-base=">, Group<m_Patmos_Group>, 
+def mpatmos_stack_base : Joined<["-"], "mpatmos-stack-base=">, Group<m_patmos_Features_Group>,
   HelpText<"Set the address for the stack cache base address symbol">, Flags<[HelpHidden]>;
-def mpatmos_shadow_stack_base : Joined<["-"], "mpatmos-shadow-stack-base=">, Group<m_Patmos_Group>, 
+def mpatmos_shadow_stack_base : Joined<["-"], "mpatmos-shadow-stack-base=">, Group<m_patmos_Features_Group>,
   HelpText<"Set the address for the shadow stack base address symbol">, Flags<[HelpHidden]>;
-def mpatmos_heap_end : Joined<["-"], "mpatmos-heap-end=">, Group<m_Patmos_Group>, 
+def mpatmos_heap_end : Joined<["-"], "mpatmos-heap-end=">, Group<m_patmos_Features_Group>,
   HelpText<"Set the address for the upper heap address symbol">, Flags<[HelpHidden]>;
 
 // Some common patmos-llc options that should be passed through. Keep in sync with patmos-llc option names and help.
 // These options might be deprecated in favor of -Xllc in the future.
-def mpatmos_disable_function_splitter : Flag<["-"], "mpatmos-disable-function-splitter">, Group<m_Patmos_llc_Group>,
+def mpatmos_disable_function_splitter : Flag<["-"], "mpatmos-disable-function-splitter">, Group<m_patmos_llc_Group>,
   HelpText<"Disable the Patmos function splitter">, Flags<[HelpHidden]>;
-def mpatmos_disable_ifcvt : Flag<["-"], "mpatmos-disable-ifcvt">, Group<m_Patmos_llc_Group>,
+def mpatmos_disable_ifcvt : Flag<["-"], "mpatmos-disable-ifcvt">, Group<m_patmos_llc_Group>,
   HelpText<"Disable if-converter for Patmos">, Flags<[HelpHidden]>;
-def mpatmos_disable_post_ra : Flag<["-"], "mpatmos-disable-post-ra">, Group<m_Patmos_llc_Group>,
+def mpatmos_disable_post_ra : Flag<["-"], "mpatmos-disable-post-ra">, Group<m_patmos_llc_Group>,
   HelpText<"Disable any post-RA scheduling">, Flags<[HelpHidden]>;
-def mpatmos_disable_stack_cache : Flag<["-"], "mpatmos-disable-stack-cache">, Group<m_Patmos_llc_Group>,
+def mpatmos_disable_stack_cache : Flag<["-"], "mpatmos-disable-stack-cache">, Group<m_patmos_llc_Group>,
   HelpText<"Disable the use of Patmos' stack cache">, Flags<[HelpHidden]>;
-def mpatmos_disable_vliw : Flag<["-"], "mpatmos-disable-vliw">, Group<m_Patmos_llc_Group>,
+def mpatmos_disable_vliw : Flag<["-"], "mpatmos-disable-vliw">, Group<m_patmos_llc_Group>,
   HelpText<"Schedule instructions only in first slot">, Flags<[HelpHidden]>;
-def mpatmos_enable_bypass_from_pml : Flag<["-"], "mpatmos-enable-bypass-from-pml">, Group<m_Patmos_llc_Group>,
+def mpatmos_enable_bypass_from_pml : Flag<["-"], "mpatmos-enable-bypass-from-pml">, Group<m_patmos_llc_Group>,
   HelpText<"Enable rewriting unanalyzable (aiT) memory accesses to bypass the cache">, Flags<[HelpHidden]>;
-def mpatmos_enable_stack_cache_analysis : Flag<["-"], "mpatmos-enable-stack-cache-analysis">, Group<m_Patmos_llc_Group>,
+def mpatmos_enable_stack_cache_analysis : Flag<["-"], "mpatmos-enable-stack-cache-analysis">, Group<m_patmos_llc_Group>,
   HelpText<"Enable the Patmos stack cache analysis">, Flags<[HelpHidden]>;
-def mpatmos_sca_preemption : Flag<["-"], "mpatmos-sca-preemption">, Group<m_Patmos_llc_Group>,
+def mpatmos_sca_preemption : Flag<["-"], "mpatmos-sca-preemption">, Group<m_patmos_llc_Group>,
   HelpText<"Enable the analysis of preemption costs during Stack Cache Analysis.">, Flags<[HelpHidden]>;
-def mpatmos_ilp_solver : Joined<["-"], "mpatmos-ilp-solver=">, Group<m_Patmos_llc_Group>,
+def mpatmos_ilp_solver : Joined<["-"], "mpatmos-ilp-solver=">, Group<m_patmos_llc_Group>,
   HelpText<"Path to an ILP solver">, Flags<[HelpHidden]>;
-def mpatmos_max_subfunction_size : Joined<["-"], "mpatmos-max-subfunction-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_max_subfunction_size : Joined<["-"], "mpatmos-max-subfunction-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Maximum size of subfunctions after function splitting, defaults to the method cache size if set to 0">, Flags<[HelpHidden]>;
-def mpatmos_method_cache_size : Joined<["-"], "mpatmos-method-cache-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_method_cache_size : Joined<["-"], "mpatmos-method-cache-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Total size of the instruction cache in bytes (default 4096)">, Flags<[HelpHidden]>;
-def mpatmos_nested_branches : Flag<["-"], "mpatmos-nested-branches">, Group<m_Patmos_llc_Group>,
+def mpatmos_nested_branches : Flag<["-"], "mpatmos-nested-branches">, Group<m_patmos_llc_Group>,
   HelpText<"Enable scheduling of branch instructions inside CFL delay slots">, Flags<[HelpHidden]>;
-def mpatmos_preferred_scc_size : Joined<["-"], "mpatmos-preferred-scc-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_preferred_scc_size : Joined<["-"], "mpatmos-preferred-scc-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Preferred maximum size for SCC subfunctions, defaults to mpatmos-preferred-subfunction-size if 0">, Flags<[HelpHidden]>;
-def mpatmos_preferred_subfunction_size : Joined<["-"], "mpatmos-preferred-subfunction-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_preferred_subfunction_size : Joined<["-"], "mpatmos-preferred-subfunction-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Preferred maximum size of subfunctions, defaults to mpatmos-max-subfunction-size if 0. Larger basic blocks and inline asm are not split">, Flags<[HelpHidden]>;
-def mpatmos_sca_remove_ensures : Flag<["-"], "mpatmos-sca-remove-ensures">, Group<m_Patmos_llc_Group>,
+def mpatmos_sca_remove_ensures : Flag<["-"], "mpatmos-sca-remove-ensures">, Group<m_patmos_llc_Group>,
   HelpText<"Remove unnecessary ensure instructions during Stack Cache Analysis">, Flags<[HelpHidden]>;
-def mpatmos_sca_serialize : Joined<["-"], "mpatmos-sca-serialize=">, Group<m_Patmos_llc_Group>,
+def mpatmos_sca_serialize : Joined<["-"], "mpatmos-sca-serialize=">, Group<m_patmos_llc_Group>,
   HelpText<"Export PML specification of generated machine code to FILE">, Flags<[HelpHidden]>;
-def mpatmos_singlepath : Joined<["-"], "mpatmos-singlepath=">, Group<m_Patmos_llc_Group>,
+def mpatmos_singlepath : Joined<["-"], "mpatmos-singlepath=">, Group<m_patmos_llc_Group>,
   HelpText<"Functions for which single-path code is generated">, Flags<[HelpHidden]>;
-def mpatmos_split_call_blocks : Flag<["-"], "mpatmos-split-call-blocks">, Group<m_Patmos_llc_Group>,
+def mpatmos_split_call_blocks : Flag<["-"], "mpatmos-split-call-blocks">, Group<m_patmos_llc_Group>,
   HelpText<"Split basic blocks containing calls into own subfunctions">, Flags<[HelpHidden]>;
-def mpatmos_stack_cache_analysis_bounds : Joined<["-"], "mpatmos-stack-cache-analysis-bounds=">, Group<m_Patmos_llc_Group>,
+def mpatmos_stack_cache_analysis_bounds : Joined<["-"], "mpatmos-stack-cache-analysis-bounds=">, Group<m_patmos_llc_Group>,
   HelpText<"File containing bounds for the stack cache analysis">, Flags<[HelpHidden]>;
-def mpatmos_stack_cache_size : Joined<["-"], "mpatmos-stack-cache-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_stack_cache_size : Joined<["-"], "mpatmos-stack-cache-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Total size of the stack cache in bytes">, Flags<[HelpHidden]>;
-def mpatmos_stack_cache_block_size : Joined<["-"], "mpatmos-stack-cache-block-size=">, Group<m_Patmos_llc_Group>,
+def mpatmos_stack_cache_block_size : Joined<["-"], "mpatmos-stack-cache-block-size=">, Group<m_patmos_llc_Group>,
   HelpText<"Block size of the stack cache in bytes">, Flags<[HelpHidden]>;
-def mpatmos_enable_block_aligned_stack_cache : Joined<["-"], "mpatmos-enable-block-aligned-stack-cache">, Group<m_Patmos_llc_Group>,
+def mpatmos_enable_block_aligned_stack_cache : Joined<["-"], "mpatmos-enable-block-aligned-stack-cache">, Group<m_patmos_llc_Group>,
   HelpText<"Enable the use of Patmos' block-aligned stack cache">, Flags<[HelpHidden]>;
 
-def mimport_pml : Joined<["-"], "mimport-pml=">, Group<m_PML_Group>, 
+def mimport_pml : Joined<["-"], "mimport-pml=">, Group<m_PML_Group>,
   HelpText<"Read external analysis results from PML file">, Flags<[HelpHidden]>;
 def mpreemit_bitcode : Joined<["-"], "mpreemit-bitcode=">, Group<m_PML_Group>,
   HelpText<"Write the final bitcode representation (before emit) to FILE">, Flags<[HelpHidden]>;
@@ -1726,22 +1731,19 @@ def nocudainc : Flag<["-"], "nocudainc">;
 def nocudalib : Flag<["-"], "nocudalib">;
 def nodefaultlibs : Flag<["-"], "nodefaultlibs">;
 def nofixprebinding : Flag<["-"], "nofixprebinding">;
-def nolibc : Flag<["-"], "nolibc">, Flags<[HelpHidden]>,
-  HelpText<"Do not link with libc.">;
+def nolibc : Flag<["-"], "nolibc">;
 def nomultidefs : Flag<["-"], "nomultidefs">;
 def nopie : Flag<["-"], "nopie">;
 def noprebind : Flag<["-"], "noprebind">;
 def noruntimelibs : Flag<["-"], "noruntimelibs">, Flags<[HelpHidden]>,
   HelpText<"Do not link with default runtime libraries.">;
 def noseglinkedit : Flag<["-"], "noseglinkedit">;
-def nostartfiles : Flag<["-"], "nostartfiles">, Flags<[HelpHidden]>,
-  HelpText<"Do not link with start files such as crt0.">;
+def nostartfiles : Flag<["-"], "nostartfiles">;
 def nostdinc : Flag<["-"], "nostdinc">;
 def nostdlibinc : Flag<["-"], "nostdlibinc">;
 def nostdincxx : Flag<["-"], "nostdinc++">, Flags<[CC1Option]>,
   HelpText<"Disable standard #include directories for the C++ standard library">;
-def nostdlib : Flag<["-"], "nostdlib">, Flags<[HelpHidden]>,
-  HelpText<"Do not link with standard libraries such as libc.">;
+def nostdlib : Flag<["-"], "nostdlib">;
 def object : Flag<["-"], "object">;
 def o : JoinedOrSeparate<["-"], "o">, Flags<[DriverOption, RenderAsInput, CC1Option, CC1AsOption]>,
   HelpText<"Write output to <file>">, MetaVarName<"<file>">;
diff --git a/include/clang/Driver/ToolChain.h b/include/clang/Driver/ToolChain.h
index 926fb3f89e..80a0d8a88e 100644
--- a/include/clang/Driver/ToolChain.h
+++ b/include/clang/Driver/ToolChain.h
@@ -225,13 +225,6 @@ public:
   /// IsBlocksDefault - Does this tool chain enable -fblocks by default.
   virtual bool IsBlocksDefault() const { return false; }
 
-  /// ShouldUseClangCompiler - Should the clang compiler be used to
-  /// handle this action.
-  virtual bool ShouldUseClangCompiler(const JobAction &JA) const;
-
-  /// IsUsingLTODefault - Does this tool chain enable -flto by default.
-  virtual bool IsUsingLTODefault() const { return false; }
-
   /// IsIntegratedAssemblerDefault - Does this tool chain enable -integrated-as
   /// by default.
   virtual bool IsIntegratedAssemblerDefault() const { return false; }
@@ -279,6 +272,9 @@ public:
   /// by default.
   virtual bool IsUnwindTablesDefault() const;
 
+  /// \brief Test whether this toolchain uses LTO by default
+  virtual bool isUsingLTODefault() const { return false; }
+
   /// \brief Test whether this toolchain defaults to PIC.
   virtual bool isPICDefault() const = 0;
 
diff --git a/include/clang/Frontend/ASTConsumers.h b/include/clang/Frontend/ASTConsumers.h
index 239f15a8a6..2ef0f3a97b 100644
--- a/include/clang/Frontend/ASTConsumers.h
+++ b/include/clang/Frontend/ASTConsumers.h
@@ -49,7 +49,7 @@ std::unique_ptr<ASTConsumer> CreateASTDeclNodeLister();
 // function declarations to stderr.
 std::unique_ptr<ASTConsumer> CreateASTViewer();
 
-ASTConsumer *CreateFlowfactExporter(StringRef filename);
+std::unique_ptr<ASTConsumer> CreateFlowfactExporter(StringRef filename);
 
 // DeclContext printer: prints out the DeclContext tree in human-readable form
 // to stderr; this is intended for debugging.
diff --git a/include/clang/Frontend/FrontendActions.h b/include/clang/Frontend/FrontendActions.h
index ce53061231..1607e6810b 100644
--- a/include/clang/Frontend/FrontendActions.h
+++ b/include/clang/Frontend/FrontendActions.h
@@ -53,8 +53,8 @@ protected:
 
 class FlowfactExportAction : public ASTFrontendAction {
 protected:
-  virtual ASTConsumer *CreateASTConsumer(CompilerInstance &CI,
-                                         StringRef InFile);
+  std::unique_ptr<ASTConsumer> CreateASTConsumer(CompilerInstance &CI,
+                                                 StringRef InFile) override;
 };
 
 class ASTDeclListAction : public ASTFrontendAction {
diff --git a/include/clang/Frontend/FrontendOptions.h b/include/clang/Frontend/FrontendOptions.h
index 644701567f..ac7e694c9c 100644
--- a/include/clang/Frontend/FrontendOptions.h
+++ b/include/clang/Frontend/FrontendOptions.h
@@ -267,6 +267,9 @@ public:
   /// \brief Auxiliary triple for CUDA compilation.
   std::string AuxTriple;
 
+  /// Name of the flowfact export file.
+  std::string FlowfactExportFile;
+
 public:
   FrontendOptions() :
     DisableFree(false), RelocatablePCH(false), ShowHelp(false),
diff --git a/include/clang/Parse/Parser.h b/include/clang/Parse/Parser.h
index 00885a5c71..d501fdc508 100644
--- a/include/clang/Parse/Parser.h
+++ b/include/clang/Parse/Parser.h
@@ -20,6 +20,7 @@
 #include "clang/Lex/CodeCompletionHandler.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/DeclSpec.h"
+#include "clang/Sema/Loopbound.h"
 #include "clang/Sema/LoopHint.h"
 #include "clang/Sema/Sema.h"
 #include "llvm/ADT/SmallVector.h"
@@ -165,6 +166,8 @@ class Parser : public CodeCompletionHandler {
   std::unique_ptr<PragmaHandler> MSSection;
   std::unique_ptr<PragmaHandler> MSRuntimeChecks;
   std::unique_ptr<PragmaHandler> OptimizeHandler;
+  std::unique_ptr<PragmaHandler> LoopboundHandler;
+  std::unique_ptr<PragmaHandler> PlatinHandler;
   std::unique_ptr<PragmaHandler> LoopHintHandler;
   std::unique_ptr<PragmaHandler> UnrollHintHandler;
   std::unique_ptr<PragmaHandler> NoUnrollHintHandler;
@@ -529,6 +532,10 @@ private:
   /// #pragma clang __debug captured
   StmtResult HandlePragmaCaptured();
 
+  /// \brief Handle the annotation token produced for
+  /// #pragma loopbound
+  void HandlePragmaLoopbound(Loopbound &LB);
+
   /// \brief Handle the annotation token produced for
   /// #pragma clang loop and #pragma unroll.
   bool HandlePragmaLoopHint(LoopHint &Hint);
@@ -1687,6 +1694,10 @@ private:
   StmtResult ParseReturnStatement();
   StmtResult ParseAsmStatement(bool &msAsm);
   StmtResult ParseMicrosoftAsmStatement(SourceLocation AsmLoc);
+  StmtResult ParsePragmaLoopbound(StmtVector &Stmts,
+                                  AllowedContsructsKind Allowed,
+                                  SourceLocation *TrailingElseLoc,
+                                  ParsedAttributesWithRange &Attrs);
   StmtResult ParsePragmaLoopHint(StmtVector &Stmts,
                                  AllowedContsructsKind Allowed,
                                  SourceLocation *TrailingElseLoc,
@@ -2497,6 +2508,8 @@ private:
   OMPClause *ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
                                       OpenMPClauseKind Kind);
 
+  /// \brief Parses declarative or executable directive.
+  StmtResult ParsePlatinPragma();
 public:
   bool ParseUnqualifiedId(CXXScopeSpec &SS, bool EnteringContext,
                           bool AllowDestructorName,
diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index 29aa642a9a..4decf2a889 100644
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -8193,6 +8193,12 @@ public:
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc);
 
+  StmtResult ActOnFlowfact(SourceLocation StartLoc,
+                           SourceLocation EndLoc,
+                           ArrayRef<int> Multipliers,
+                           ArrayRef<std::string> Markers,
+                           int Rhs);
+
   /// \brief The kind of conversion being performed.
   enum CheckedConversionKind {
     /// \brief An implicit conversion.
diff --git a/lib/AST/Stmt.cpp b/lib/AST/Stmt.cpp
index 4c79f5bd1f..afbaced720 100644
--- a/lib/AST/Stmt.cpp
+++ b/lib/AST/Stmt.cpp
@@ -1109,3 +1109,35 @@ bool CapturedStmt::capturesVariable(const VarDecl *Var) const {
 
   return false;
 }
+
+Flowfact::Flowfact(SourceRange Range,
+                   ArrayRef<int> Multipliers,
+                   ArrayRef<std::string> Markers,
+                   int rhs)
+  : Stmt(FlowfactClass), Range(Range),
+    Multipliers(Multipliers.begin(), Multipliers.end()),
+    Markers(Markers.begin(), Markers.end()),
+    RHS(rhs) {
+    assert(Markers.size() == Multipliers.size());
+}
+
+Flowfact *Flowfact::Create(const ASTContext &C,
+                           SourceRange Range,
+                           ArrayRef<int> Multipliers,
+                           ArrayRef<std::string> Markers,
+                           int rhs) {
+  // XXX this is not he clang way of allocating the Flowfact statement, at least
+  // the LHS terms should be stored as sub-expressions and allocated here
+  void *Mem = C.Allocate(sizeof(Flowfact),
+                         llvm::alignOf<Flowfact>());
+  return new (Mem) Flowfact(Range, Multipliers, Markers, rhs);
+}
+
+Flowfact *Flowfact::CreateEmpty(const ASTContext &C) {
+  void *Mem = C.Allocate(sizeof(Flowfact),
+                         llvm::alignOf<Flowfact>());
+  return new (Mem) Flowfact(SourceRange(),
+                            ArrayRef<int>(),
+                            ArrayRef<std::string>(),
+                            -1);
+}
diff --git a/lib/AST/StmtPrinter.cpp b/lib/AST/StmtPrinter.cpp
index 69f52f52b6..f02ecfe05d 100644
--- a/lib/AST/StmtPrinter.cpp
+++ b/lib/AST/StmtPrinter.cpp
@@ -581,6 +581,15 @@ void StmtPrinter::VisitSEHLeaveStmt(SEHLeaveStmt *Node) {
   if (Policy.IncludeNewlines) OS << "\n";
 }
 
+void StmtPrinter::VisitFlowfact(Flowfact *Node) {
+  unsigned i;
+  Indent() << "Flowfact (";
+  for (i = 0; i < Node->getNumLhsTerms() - 1; ++i)
+    OS << Node->Multipliers[i] << " " << Node->Markers[i] << " + ";
+  OS << Node->Multipliers[i] << " " << Node->Markers[i]
+    << " <= " << Node->RHS << ")\n";
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenMP clauses printing methods
 //===----------------------------------------------------------------------===//
diff --git a/lib/AST/StmtProfile.cpp b/lib/AST/StmtProfile.cpp
index 175a43abbf..3209eb1792 100644
--- a/lib/AST/StmtProfile.cpp
+++ b/lib/AST/StmtProfile.cpp
@@ -611,6 +611,10 @@ void StmtProfiler::VisitOMPDistributeDirective(
   VisitOMPLoopDirective(S);
 }
 
+void StmtProfiler::VisitFlowfact(const Flowfact *Node) {
+  VisitStmt(Node);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }
diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index af8aea0929..ae769caaf9 100644
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -5836,6 +5836,154 @@ const Builtin::Info HexagonTargetInfo::BuiltinInfo[] = {
 #include "clang/Basic/BuiltinsHexagon.def"
 };
 
+
+// Patmos abstract base class
+// TODO: builtins
+class PatmosTargetInfo : public TargetInfo {
+  bool SoftFloat : 1;
+public:
+  PatmosTargetInfo(const llvm::Triple &triple) : TargetInfo(triple)  {
+    BigEndian = true;
+    SoftFloat = true;
+    // Keep in sync with PatmosTargetMachine and compiler-rt/lib/patmos/*.ll
+    DataLayoutString =
+      "E-S32-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:32-f32:32:32-f64:32:32-a0:0:32-v64:32:32-v128:32:32-n32";
+    // Note: those values must be kept in sync with the DescriptionString!
+    DoubleAlign = 32;
+    LongLongAlign = 32;
+    LongDoubleAlign = 32;
+    SuitableAlign = 32;
+    PreferWidthAligned = false;
+    // Keep {|} as they are in inline asm
+    NoAsmVariants = true;
+  }
+
+  void setFeatureEnabled(llvm::StringMap<bool> &Features,
+                         StringRef Name,
+                         bool Enabled) const override {
+    if (Name == "hard-float" || Name == "soft-float") {
+      Features[Name] = Enabled;
+    }
+  }
+
+  bool handleTargetFeatures(std::vector<std::string> &Features,
+                            DiagnosticsEngine &Diags) override {
+    SoftFloat = true;
+    for (unsigned i = 0, e = Features.size(); i != e; ++i) {
+      if (Features[i] == "+hard-float")
+        SoftFloat = false;
+      if (Features[i] == "+soft-float")
+        SoftFloat = true;
+    }
+    return true;
+  }
+
+  bool hasFeature(StringRef Feature) const override {
+    return llvm::StringSwitch<bool>(Feature)
+             .Case("softfloat", SoftFloat)
+             .Case("patmos", true)
+             .Default(false);
+  }
+
+  void getTargetDefines(const LangOptions &Opts,
+                                MacroBuilder &Builder) const {
+    // Target identification.
+    Builder.defineMacro("__patmos__");
+    Builder.defineMacro("__PATMOS__");
+
+    if (SoftFloat)
+      Builder.defineMacro("SOFT_FLOAT", "1");
+  }
+
+  ArrayRef<Builtin::Info> getTargetBuiltins() const override {
+    return None;
+  }
+
+  const char *getClobbers() const override {
+    return "";
+  }
+
+  BuiltinVaListKind getBuiltinVaListKind() const override {
+    return TargetInfo::CharPtrBuiltinVaList;
+  }
+
+  ArrayRef<const char *> getGCCRegNames() const override
+  {
+    // TODO can't we get this from the backend tables somehow??
+    static const char * const GCCRegNames[] = {
+      // CPU register names
+      // Must match second column of GCCRegAliases
+      // The names here must match the register enum names in the .td file,
+      // not the register name string value (case insensitive).
+      "$r0",   "$r1",   "$r2",   "$r3",   "$r4",   "$r5",   "$r6",   "$r7",
+      "$r8",   "$r9",   "$r10",  "$r11",  "$r12",  "$r13",  "$r14",  "$r15",
+      "$r16",  "$r17",  "$r18",  "$r19",  "$r20",  "$r21",  "$r22",  "$r23",
+      "$r24",  "$r25",  "$r26",  "$rtr",  "$rfp",  "$rsp",  "$rfb",  "$rfo",
+      // Predicates
+      "$p0",  "$p1",  "$p2",  "$p3",  "$p4",  "$p5",  "$p6",  "$p7",
+      // Special registers
+      "$s0",  "$sm",  "$sl",  "$sh",  "$s4",  "$ss",  "$st",  "$s7",
+      "$s8",  "$s9",  "$s10", "$s11", "$s12", "$s13", "$s14", "$s15"
+    };
+    return llvm::makeArrayRef(GCCRegNames);
+  }
+
+  ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
+    static const GCCRegAlias GCCRegAliases[] = {
+        { { "$r27" }, "$rtr" },
+        { { "$r28" }, "$rfp" },
+        { { "$r29" }, "$rsp" },
+        { { "$r30" }, "$rfb" },
+        { { "$r31" }, "$rfo" },
+        { { "$s1"  }, "$sm"  },
+        { { "$s2"  }, "$sl"  },
+        { { "$s3"  }, "$sh"  },
+        { { "$s5"  }, "$ss"  },
+        { { "$s6"  }, "$st"  }
+    };
+    return llvm::makeArrayRef(GCCRegAliases);
+  }
+
+  bool validateAsmConstraint(const char *&Name,
+                             TargetInfo::ConstraintInfo &Info) const override
+  {
+    // clang actually accepts a few generic register constraints more (i,n,m,o,g,..),
+    // not much we can do about it.. For completeness, we list all currently supported
+    // constraints here.
+
+    switch (*Name) {
+    default:
+      return false;
+
+    case 'r': // CPU registers.
+    // TODO do not accept read-only or special registers here
+    case 'R': // r0-r31, currently same as 'r'
+    case 'S': // sz-s15
+    case 'P': // p0-p7
+    // TODO define more classes for subsets of registers (r10-r28, ..)?
+      Info.setAllowsRegister();
+      return true;
+    case '{':
+      Name++;
+      if (!*Name || *Name != '$')
+        return false;
+      Name++;
+      while (*Name) {
+        if (*Name == '}') {
+          return true;
+        }
+        if (*Name != 'r' && *Name != 's' && *Name != 'p' &&
+            (*Name < '0' || *Name > '9')) {
+          return false;
+        }
+        Name++;
+      }
+      return false;
+    }
+  }
+};
+
+
 // Shared base class for SPARC v8 (32-bit) and SPARC v9 (64-bit).
 class SparcTargetInfo : public TargetInfo {
   static const TargetInfo::GCCRegAlias GCCRegAliases[];
@@ -7528,6 +7676,9 @@ static TargetInfo *AllocateTarget(const llvm::Triple &Triple) {
   case llvm::Triple::hexagon:
     return new HexagonTargetInfo(Triple);
 
+  case llvm::Triple::patmos:
+    return new PatmosTargetInfo(Triple);
+
   case llvm::Triple::aarch64:
     if (Triple.isOSDarwin())
       return new DarwinAArch64TargetInfo(Triple);
diff --git a/lib/CodeGen/BackendUtil.cpp b/lib/CodeGen/BackendUtil.cpp
index afae8de175..6d746c25ee 100644
--- a/lib/CodeGen/BackendUtil.cpp
+++ b/lib/CodeGen/BackendUtil.cpp
@@ -407,12 +407,6 @@ void EmitAssemblyHelper::CreatePasses(FunctionInfoIndex *FunctionIndex) {
     FPM->add(createVerifierPass());
   PMBuilder.populateFunctionPassManager(*FPM);
 
-  // Add baseline opts when the build is going to optimize bitcode later
-  if (CodeGenOpts.EnableLLVMBaselineOpts) {
-    assert(CodeGenOpts.DisableLLVMOpts && "baseline opts imply DisableLLVMOpts");
-    PMBuilder.populateFPMBaseline(*FPM);
-  }
-
   // Set up the per-module pass manager.
   if (!CodeGenOpts.RewriteMapFiles.empty())
     addSymbolRewriterPass(CodeGenOpts, MPM);
diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp
index 3a4fdad1af..f1ac41d245 100644
--- a/lib/CodeGen/CGBuiltin.cpp
+++ b/lib/CodeGen/CGBuiltin.cpp
@@ -64,7 +64,7 @@ llvm::Value *CodeGenModule::getLLVMIntrinsicFunction(const FunctionDecl *FD,
     cast<llvm::FunctionType>(getTypes().ConvertType(FD->getType()));
   switch (BuiltinID) {
   default: // not yet supported
-    this->ErrorUnsupported(FD, Context.BuiltinInfo.GetName(BuiltinID));
+    this->ErrorUnsupported(FD, Context.BuiltinInfo.getName(BuiltinID));
     return llvm::UndefValue::get(Ty);
     ;
   case Builtin::BI__llvm_pcmarker:
@@ -1994,6 +1994,12 @@ RValue CodeGenFunction::EmitBuiltinExpr(const FunctionDecl *FD,
   if (getContext().BuiltinInfo.isPredefinedLibFunction(BuiltinID))
     return emitLibraryCall(*this, FD, E, EmitScalarExpr(E->getCallee()));
 
+  // If this a call to a LLVM intrinsic (e.g. __llvm_pcmarker), generate the
+  // corresponding call to the intrinsic function (e.g., @llvm.pcmarker)
+  if (getContext().BuiltinInfo.isLLVMIntrinsicFunction(BuiltinID))
+    return emitLibraryCall(*this, FD, E,
+                           CGM.getLLVMIntrinsicFunction(FD, BuiltinID));
+
   // Check that a call to a target specific builtin has the correct target
   // features.
   // This is down here to avoid non-target specific builtins, however, if
diff --git a/lib/CodeGen/CGStmt.cpp b/lib/CodeGen/CGStmt.cpp
index 1810bd1a71..f08e3aaa08 100644
--- a/lib/CodeGen/CGStmt.cpp
+++ b/lib/CodeGen/CGStmt.cpp
@@ -181,6 +181,8 @@ void CodeGenFunction::EmitStmt(const Stmt *S) {
   case Stmt::SEHTryStmtClass:
     EmitSEHTryStmt(cast<SEHTryStmt>(*S));
     break;
+  case Stmt::FlowfactClass:
+    // flowfacts are not being translated at this point
   case Stmt::OMPParallelDirectiveClass:
     EmitOMPParallelDirective(cast<OMPParallelDirective>(*S));
     break;
@@ -610,6 +612,78 @@ void CodeGenFunction::EmitIfStmt(const IfStmt &S) {
   EmitBlock(ContBlock, true);
 }
 
+// This inserts a loopbound intrinsic into the basic block that
+// is the header of an emitted loop.
+// We assume here that the block is already filled with instructions.
+void CodeGenFunction::EmitHeaderBounds(llvm::BasicBlock *Header,
+                                       const ArrayRef<const Attr *> &Attrs) {
+  // Return if there are no hints.
+  if (Attrs.empty())
+    return;
+
+  for (unsigned i = 0; i < Attrs.size(); ++i) {
+    const Attr *A = Attrs[i];
+    const LoopboundAttr *LB = dyn_cast<LoopboundAttr>(A);
+
+    // Skip non loopbound attributes
+    if (!LB) continue;
+
+    SmallVector<llvm::Value*, 16> Args;
+    llvm::Value *MinVal = llvm::ConstantInt::get(Int32Ty, LB->getMin());
+    llvm::Value *MaxVal = llvm::ConstantInt::get(Int32Ty, LB->getMax());
+
+    Args.push_back(MinVal);
+    Args.push_back(MaxVal);
+    llvm::BasicBlock::iterator I = Header->getFirstInsertionPt();
+
+    llvm::Value *Callee =
+      CGM.getIntrinsic(llvm::Intrinsic::loopbound);
+    llvm::CallSite CS = llvm::CallInst::Create(Callee, Args, "", &*I);
+    (void) CS;
+  }
+
+}
+
+void CodeGenFunction::EmitCondBrBounds(llvm::LLVMContext &Context,
+                                       llvm::BranchInst *CondBr,
+                                       const ArrayRef<const Attr *> &Attrs) {
+  // Return if there are no hints.
+  if (Attrs.empty())
+    return;
+
+  // Add loopbounds to the metadata on the conditional branch.
+  SmallVector<llvm::Metadata *, 2> Metadata;
+  for (unsigned i = 0; i < Attrs.size(); ++i) {
+    const Attr *A = Attrs[i];
+    const LoopboundAttr *LB = dyn_cast<LoopboundAttr>(A);
+
+    // Skip non loopbound attributes
+    if (!LB) continue;
+
+    const char *MetadataName = "llvm.loop.bound";
+    llvm::MDString *Name = llvm::MDString::get(Context, MetadataName);
+    llvm::Value *MinVal = llvm::ConstantInt::get(Int32Ty, LB->getMin());
+    llvm::Value *MaxVal = llvm::ConstantInt::get(Int32Ty, LB->getMax());
+
+    SmallVector<llvm::Metadata *, 3> OpValues;
+    OpValues.push_back(Name);
+    OpValues.push_back(llvm::ValueAsMetadata::get(MinVal));
+    OpValues.push_back(llvm::ValueAsMetadata::get(MaxVal));
+
+    // Set or overwrite metadata indicated by Name.
+    Metadata.push_back(llvm::MDNode::get(Context, OpValues));
+  }
+
+  if (!Metadata.empty()) {
+    // Add llvm.loop MDNode to CondBr.
+    llvm::MDNode *LoopID = llvm::MDNode::get(Context, Metadata);
+    LoopID->replaceOperandWith(0, LoopID); // First op points to itself.
+
+    CondBr->setMetadata("llvm.loop", LoopID);
+  }
+
+}
+
 void CodeGenFunction::EmitWhileStmt(const WhileStmt &S,
                                     ArrayRef<const Attr *> WhileAttrs) {
   // Emit the header for the loop, which will also become
@@ -656,7 +730,9 @@ void CodeGenFunction::EmitWhileStmt(const WhileStmt &S,
     llvm::BasicBlock *ExitBlock = LoopExit.getBlock();
     if (ConditionScope.requiresCleanups())
       ExitBlock = createBasicBlock("while.exit");
-    Builder.CreateCondBr(
+
+    llvm::BranchInst *CondBr =
+      Builder.CreateCondBr(
         BoolCondVal, LoopBody, ExitBlock,
         createProfileWeightsForLoop(S.getCond(), getProfileCount(S.getBody())));
 
@@ -664,6 +740,9 @@ void CodeGenFunction::EmitWhileStmt(const WhileStmt &S,
       EmitBlock(ExitBlock);
       EmitBranchThroughCleanup(LoopExit);
     }
+
+    // Attach metadata to loop body conditional branch.
+    EmitCondBrBounds(LoopBody->getContext(), CondBr, WhileAttrs);
   }
 
   // Emit the loop body.  We have to emit this in a cleanup scope
@@ -676,7 +755,7 @@ void CodeGenFunction::EmitWhileStmt(const WhileStmt &S,
   }
 
   // Insert loopbound instrinsic
-  EmitHeaderBounds(LoopHeader.getBlock(), Attrs);
+  EmitHeaderBounds(LoopHeader.getBlock(), WhileAttrs);
 
 
   BreakContinueStack.pop_back();
@@ -720,6 +799,9 @@ void CodeGenFunction::EmitDoStmt(const DoStmt &S,
     EmitStmt(S.getBody());
   }
 
+  // Insert loopbound instrinsic
+  EmitHeaderBounds(LoopBody, DoAttrs);
+
   EmitBlock(LoopCond.getBlock());
 
   // C99 6.8.5.2: "The evaluation of the controlling expression takes place
@@ -742,9 +824,13 @@ void CodeGenFunction::EmitDoStmt(const DoStmt &S,
   // As long as the condition is true, iterate the loop.
   if (EmitBoolCondBranch) {
     uint64_t BackedgeCount = getProfileCount(S.getBody()) - ParentCount;
-    Builder.CreateCondBr(
+    llvm::BranchInst *CondBr =
+      Builder.CreateCondBr(
         BoolCondVal, LoopBody, LoopExit.getBlock(),
         createProfileWeightsForLoop(S.getCond(), BackedgeCount));
+
+    // Attach metadata to loop body conditional branch.
+    EmitCondBrBounds(LoopBody->getContext(), CondBr, DoAttrs);
   }
 
   LoopStack.pop();
@@ -809,10 +895,14 @@ void CodeGenFunction::EmitForStmt(const ForStmt &S,
     // C99 6.8.5p2/p4: The first substatement is executed if the expression
     // compares unequal to 0.  The condition must be a scalar type.
     llvm::Value *BoolCondVal = EvaluateExprAsBool(S.getCond());
-    Builder.CreateCondBr(
+    llvm::BranchInst *CondBr =
+      Builder.CreateCondBr(
         BoolCondVal, ForBody, ExitBlock,
         createProfileWeightsForLoop(S.getCond(), getProfileCount(S.getBody())));
 
+    // Attach metadata to loop body conditional branch.
+    EmitCondBrBounds(ForBody->getContext(), CondBr, ForAttrs);
+
     if (ExitBlock != LoopExit.getBlock()) {
       EmitBlock(ExitBlock);
       EmitBranchThroughCleanup(LoopExit);
@@ -846,7 +936,7 @@ void CodeGenFunction::EmitForStmt(const ForStmt &S,
   EmitBranch(CondBlock);
 
   // Insert loopbound instrinsic
-  EmitHeaderBounds(CondBlock, Attrs);
+  EmitHeaderBounds(CondBlock, ForAttrs);
 
 
   ForScope.ForceCleanup();
@@ -888,10 +978,14 @@ CodeGenFunction::EmitCXXForRangeStmt(const CXXForRangeStmt &S,
   // The body is executed if the expression, contextually converted
   // to bool, is true.
   llvm::Value *BoolCondVal = EvaluateExprAsBool(S.getCond());
-  Builder.CreateCondBr(
+  llvm::BranchInst *CondBr =
+    Builder.CreateCondBr(
       BoolCondVal, ForBody, ExitBlock,
       createProfileWeightsForLoop(S.getCond(), getProfileCount(S.getBody())));
 
+  // Attach metadata to loop body conditional branch.
+  EmitCondBrBounds(ForBody->getContext(), CondBr, ForAttrs);
+
   if (ExitBlock != LoopExit.getBlock()) {
     EmitBlock(ExitBlock);
     EmitBranchThroughCleanup(LoopExit);
@@ -923,7 +1017,7 @@ CodeGenFunction::EmitCXXForRangeStmt(const CXXForRangeStmt &S,
   EmitBranch(CondBlock);
 
   // Insert loopbound instrinsic
-  EmitHeaderBounds(CondBlock, Attrs);
+  EmitHeaderBounds(CondBlock, ForAttrs);
 
   ForScope.ForceCleanup();
 
diff --git a/lib/CodeGen/CodeGenFunction.h b/lib/CodeGen/CodeGenFunction.h
index 3e9016adb3..21fc3bd77c 100644
--- a/lib/CodeGen/CodeGenFunction.h
+++ b/lib/CodeGen/CodeGenFunction.h
@@ -1485,9 +1485,7 @@ public:
   llvm::BasicBlock *createBasicBlock(const Twine &name = "",
                                      llvm::Function *parent = nullptr,
                                      llvm::BasicBlock *before = nullptr) {
-#ifdef NDEBUG
-    return llvm::BasicBlock::Create(getLLVMContext(), "", parent, before);
-#else
+    // Patmos-specific: always keep basic block names for PML export
     return llvm::BasicBlock::Create(getLLVMContext(), name, parent, before);
   }
 
@@ -2139,6 +2137,11 @@ public:
   void EmitIndirectGotoStmt(const IndirectGotoStmt &S);
   void EmitIfStmt(const IfStmt &S);
 
+  void EmitCondBrBounds(llvm::LLVMContext &Context, llvm::BranchInst *CondBr,
+                        const ArrayRef<const Attr *> &Attrs);
+  void EmitHeaderBounds(llvm::BasicBlock *Header,
+                        const ArrayRef<const Attr *> &Attrs);
+
   void EmitWhileStmt(const WhileStmt &S,
                      ArrayRef<const Attr *> Attrs = None);
   void EmitDoStmt(const DoStmt &S, ArrayRef<const Attr *> Attrs = None);
diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 750a4fb29b..6ce735db93 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -6974,6 +6974,33 @@ Address XCoreABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr,
   return Val;
 }
 
+//===----------------------------------------------------------------------===//
+// Patmos ABI Implementation
+//===----------------------------------------------------------------------===//
+namespace {
+class PatmosABIInfo : public DefaultABIInfo {
+public:
+  PatmosABIInfo(CodeGen::CodeGenTypes &CGT) : DefaultABIInfo(CGT) {}
+};
+
+class PatmosTargetCodeGenInfo : public TargetCodeGenInfo {
+public:
+  PatmosTargetCodeGenInfo(CodeGenTypes &CGT)
+    : TargetCodeGenInfo(new PatmosABIInfo(CGT)) {}
+
+  void SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
+                           CodeGen::CodeGenModule &CGM) const {
+    const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
+    if (!FD) return;
+    llvm::Function *Fn = cast<llvm::Function>(GV);
+    if (FD->hasAttr<SinglePathAttr>()) {
+      Fn->addFnAttr("sp-root");
+      Fn->addFnAttr(llvm::Attribute::NoInline);
+    }
+  }
+};
+}
+
 /// During the expansion of a RecordType, an incomplete TypeString is placed
 /// into the cache as a means to identify and break recursion.
 /// If there is a Recursive encoding in the cache, it is swapped out and will
diff --git a/lib/Driver/CMakeLists.txt b/lib/Driver/CMakeLists.txt
index fa0430e211..f0733c384c 100644
--- a/lib/Driver/CMakeLists.txt
+++ b/lib/Driver/CMakeLists.txt
@@ -1,4 +1,5 @@
 set(LLVM_LINK_COMPONENTS
+  Object
   Option
   Support
   )
diff --git a/lib/Driver/Driver.cpp b/lib/Driver/Driver.cpp
index 1e0a48d529..19a57edec3 100644
--- a/lib/Driver/Driver.cpp
+++ b/lib/Driver/Driver.cpp
@@ -370,8 +370,8 @@ static llvm::Triple computeTargetTriple(StringRef DefaultTargetTriple,
 
 // \brief Parse the LTO options and record the type of LTO compilation
 // based on which -f(no-)?lto(=.*)? option occurs last.
-void Driver::setLTOMode(const llvm::opt::ArgList &Args) {
-  LTOMode = LTOK_None;
+void Driver::setLTOMode(const ToolChain &TC, const llvm::opt::ArgList &Args) {
+  LTOMode = TC.isUsingLTODefault() ? LTOK_Full : LTOK_None;
   if (!Args.hasFlag(options::OPT_flto, options::OPT_flto_EQ,
                     options::OPT_fno_lto, false))
     return;
@@ -477,8 +477,6 @@ Compilation *Driver::BuildCompilation(ArrayRef<const char *> ArgList) {
                     .Default(SaveTempsCwd);
   }
 
-  setLTOMode(Args);
-
   std::unique_ptr<llvm::opt::InputArgList> UArgs =
       llvm::make_unique<InputArgList>(std::move(Args));
 
@@ -489,6 +487,8 @@ Compilation *Driver::BuildCompilation(ArrayRef<const char *> ArgList) {
   const ToolChain &TC =
       getToolChain(*UArgs, computeTargetTriple(DefaultTargetTriple, *UArgs));
 
+  setLTOMode(TC, *UArgs);
+
   // The compilation takes ownership of Args.
   Compilation *C = new Compilation(*this, TC, UArgs.release(), TranslatedArgs);
 
@@ -2290,6 +2290,9 @@ const ToolChain &Driver::getToolChain(const ArgList &Args,
       case llvm::Triple::hexagon:
         TC = new toolchains::HexagonToolChain(*this, Target, Args);
         break;
+      case llvm::Triple::patmos:
+        TC = new toolchains::PatmosToolChain(*this, Target, Args);
+        break;
       case llvm::Triple::xcore:
         TC = new toolchains::XCoreToolChain(*this, Target, Args);
         break;
diff --git a/lib/Driver/ToolChains.cpp b/lib/Driver/ToolChains.cpp
index 99c7b8e68c..784a94c9ad 100644
--- a/lib/Driver/ToolChains.cpp
+++ b/lib/Driver/ToolChains.cpp
@@ -2726,6 +2726,133 @@ const StringRef HexagonToolChain::GetTargetCPUVersion(const ArgList &Args) {
 }
 // End Hexagon
 
+/// PatmosToolChain - A tool chain using the llvm bitcode tools to perform
+/// all sub-commands.
+
+PatmosToolChain::PatmosToolChain(const Driver &D, const llvm::Triple& Triple,
+                                const ArgList &Args)
+  : ToolChain(D, Triple, Args) {
+  // Get install path to find tools and libraries
+  std::string Path(D.getInstalledDir());
+
+  // tools?
+  getProgramPaths().push_back(Path);
+  if (llvm::sys::fs::exists(Path + "/bin/"))
+    getProgramPaths().push_back(Path + "/bin/");
+  if (llvm::sys::fs::exists(Path + "/../bin/"))
+    getProgramPaths().push_back(Path + "/../bin/");
+
+  // add lib to search paths so that we can look for LLVMgold.so
+  if (llvm::sys::fs::exists(Path + "/lib/"))
+    getProgramPaths().push_back(Path + "/lib/");
+  if (llvm::sys::fs::exists(Path + "/../lib/"))
+    getProgramPaths().push_back(Path + "/../lib/");
+
+  // TODO merge with ComputeLLVMTriple below somehow?
+  std::string TripleString = Triple.getTriple();
+  if (TripleString == "patmos")
+    TripleString = "patmos-unknown-unknown-elf";
+
+  // newlib libraries and includes?
+  // checking patmos-unknown-elf for backward-compatibility reasons
+  if (llvm::sys::fs::exists(Path + "/" + TripleString + "/"))
+    getFilePaths().push_back(Path + "/" + TripleString + "/");
+  else if (llvm::sys::fs::exists(Path + "/patmos-unknown-elf/"))
+    getFilePaths().push_back(Path + "/patmos-unknown-elf/");
+
+  if (llvm::sys::fs::exists(Path + "/../" + TripleString + "/"))
+    getFilePaths().push_back(Path + "/../" + TripleString + "/");
+  else if (llvm::sys::fs::exists(Path + "/../patmos-unknown-elf/"))
+    getFilePaths().push_back(Path + "/../patmos-unknown-elf/");
+}
+
+PatmosToolChain::~PatmosToolChain() {
+}
+
+std::string PatmosToolChain::ComputeLLVMTriple(const ArgList &Args,
+                               types::ID InputType) const
+{
+  if (getTriple().getArch() != llvm::Triple::patmos) {
+    llvm_unreachable("Invalid architecture for Patmos tool chain");
+  }
+
+  std::string Triple = getTripleString();
+
+  // This is a bit of a workaround: when we call patmos-clang without
+  // -target, then clang uses 'patmos' as default target (the prefix of the
+  // program call). To avoid target-name mismatches, we normalize that to
+  // the full default triple.
+  if (Triple == "patmos") {
+    return "patmos-unknown-unknown-elf";
+  }
+
+  return Triple;
+}
+
+bool PatmosToolChain::IsMathErrnoDefault() const {
+  return true;
+}
+
+bool PatmosToolChain::IsUnwindTablesDefault() const {
+  return false;
+}
+
+const char *PatmosToolChain::GetDefaultRelocationModel() const {
+  return "static";
+}
+
+const char *PatmosToolChain::GetForcedPicModel() const {
+  return 0;
+}
+
+void PatmosToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs,
+                                                ArgStringList &CC1Args) const
+{
+  if (!DriverArgs.hasArg(options::OPT_nostdinc) &&
+      !DriverArgs.hasArg(options::OPT_nostdlibinc)) {
+    const ToolChain::path_list &filePaths = getFilePaths();
+    for(ToolChain::path_list::const_iterator i = filePaths.begin(),
+        ie = filePaths.end(); i != ie; i++) {
+      // construct a library search path
+      CC1Args.push_back("-isystem");
+      CC1Args.push_back(DriverArgs.MakeArgString(Twine(*i) + "include/"));
+    }
+  }
+}
+
+Tool *PatmosToolChain::getTool(Action::ActionClass AC) const
+{
+  if (AC == Action::CompileJobClass ||
+      AC == Action::BackendJobClass) {
+    // Use a special clang driver that supports compiling to ELF with
+    // -fpatmos-emit-obj
+    return getPatmosClang();
+  }
+  return ToolChain::getTool(AC);
+}
+
+Tool *PatmosToolChain::SelectTool(const JobAction &JA) const {
+  Action::ActionClass AC = JA.getKind();
+  if (AC == Action::CompileJobClass ||
+      AC == Action::BackendJobClass) {
+    // override the default tool selection for the compiler
+    return getTool(AC);
+  }
+  return ToolChain::SelectTool(JA);
+}
+
+Tool *PatmosToolChain::getPatmosClang() const {
+  if (!PatmosClang)
+    PatmosClang.reset(new tools::patmos::Compile(*this));
+  return PatmosClang.get();
+}
+
+Tool* PatmosToolChain::buildLinker() const {
+  return new tools::patmos::Link(*this);
+}
+// End Patmos
+
+
 /// AMDGPU Toolchain
 AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
                                  const ArgList &Args)
diff --git a/lib/Driver/ToolChains.h b/lib/Driver/ToolChains.h
index f940e5847e..13e2977f97 100644
--- a/lib/Driver/ToolChains.h
+++ b/lib/Driver/ToolChains.h
@@ -893,6 +893,43 @@ public:
       const llvm::opt::ArgList &Args);
 };
 
+class LLVM_LIBRARY_VISIBILITY PatmosToolChain : public ToolChain {
+private:
+  mutable std::unique_ptr<Tool> PatmosClang;
+
+  Tool *getPatmosClang() const;
+public:
+  PatmosToolChain(const Driver &D, const llvm::Triple& Triple, 
+                 const llvm::opt::ArgList& Args);
+  ~PatmosToolChain();
+
+  virtual std::string ComputeLLVMTriple(const llvm::opt::ArgList &Args,
+                                 types::ID InputType = types::TY_INVALID) const;
+
+  Tool *SelectTool(const JobAction &JA) const override;
+
+  bool IsMathErrnoDefault() const;
+  bool IsUnwindTablesDefault() const;
+  const char* GetDefaultRelocationModel() const;
+  const char* GetForcedPicModel() const;
+
+  virtual void AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
+                                         llvm::opt::ArgStringList &CC1Args) const;
+
+  virtual bool SupportsProfiling() const { return false; }
+  virtual bool IsIntegratedAssemblerDefault() const { return true; }
+
+  virtual bool isUsingLTODefault() const { return true; }
+
+  virtual bool isPICDefault() const { return false; }
+  virtual bool isPIEDefault() const { return false; }
+  virtual bool isPICDefaultForced() const { return false; }
+protected:
+  Tool *getTool(Action::ActionClass AC) const override;
+
+  Tool *buildLinker() const override;
+};
+
 class LLVM_LIBRARY_VISIBILITY AMDGPUToolChain : public Generic_ELF {
 protected:
   Tool *buildLinker() const override;
diff --git a/lib/Driver/Tools.cpp b/lib/Driver/Tools.cpp
index edbba26c34..36400db33e 100644
--- a/lib/Driver/Tools.cpp
+++ b/lib/Driver/Tools.cpp
@@ -29,6 +29,7 @@
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/Twine.h"
+#include "llvm/IR/LLVMContext.h"
 #include "llvm/Config/config.h"
 #include "llvm/Object/Archive.h"
 #include "llvm/Option/Arg.h"
@@ -2072,6 +2073,98 @@ void Clang::AddHexagonTargetArgs(const ArgList &Args,
   CmdArgs.push_back("-machine-sink-split=0");
 }
 
+static StringRef getPatmosFloatABI(const Driver &D, const ArgList &Args,
+                                   const llvm::Triple &Triple,
+                                   bool &DefaultChanged) {
+  // Select the float ABI as determined by -msoft-float, -mhard-float,
+  // and -mfloat-abi=.
+
+  // TODO determine default FloatABI based on the processor subtarget features
+  StringRef FloatABI = "soft";
+  // Set to true when the user selects something different from the (subtarget) default.
+  DefaultChanged = false;
+
+  if (Arg *A = Args.getLastArg(options::OPT_msoft_float,
+                               options::OPT_mhard_float,
+                               options::OPT_mfloat_abi_EQ,
+                               options::OPT_mno_soft_float)) {
+    if (A->getOption().matches(options::OPT_msoft_float)) {
+      FloatABI = "soft";
+    } else if (A->getOption().matches(options::OPT_mhard_float)) {
+      FloatABI = "hard";
+      DefaultChanged = true;
+    } else if (A->getOption().matches(options::OPT_mno_soft_float)) {
+      FloatABI = "none";
+      DefaultChanged = true;
+    } else {
+      FloatABI = A->getValue();
+      if (FloatABI != "soft" && FloatABI != "hard" &&
+          FloatABI != "none" && FloatABI != "simple") {
+        D.Diag(diag::err_drv_invalid_mfloat_abi) << A->getAsString(Args);
+        FloatABI = "soft";
+      }
+      DefaultChanged = (FloatABI != "soft");
+    }
+  }
+ 
+  return FloatABI;
+}
+
+void Clang::AddPatmosTargetArgs(const ArgList &Args,
+                                ArgStringList &CmdArgs) const {
+
+  // we do not want to have the host includes here
+  CmdArgs.push_back("-nostdsysteminc");
+
+  // Set correct floating-point flags
+  bool Changed;
+  StringRef FloatABI = getPatmosFloatABI(getToolChain().getDriver(), Args,
+                                         getToolChain().getTriple(), Changed);
+
+  if (FloatABI == "soft") {
+    // Floating point operations and argument passing are soft.
+    CmdArgs.push_back("-msoft-float");
+    CmdArgs.push_back("-mfloat-abi");
+    CmdArgs.push_back("soft");
+  } else if (FloatABI == "none") {
+    // Do not use floats at all
+    // TODO anything else to add?
+    CmdArgs.push_back("-mfloat-abi");
+    CmdArgs.push_back("hard");
+  }
+  else {
+    // Floating point operations and argument passing are hard.
+    if (FloatABI != "hard") {
+      getToolChain().getDriver().Diag(diag::err_drv_invalid_mfloat_abi) << FloatABI;
+    }
+    CmdArgs.push_back("-mfloat-abi");
+    CmdArgs.push_back("hard");
+
+    // Pass float mode to PatmosTargetInfo
+    CmdArgs.push_back("-target-feature");
+    CmdArgs.push_back("+hard-float");
+  }
+
+  // Disable nested loop separation for loops with multiple backedges
+  // during Natural Loop Canonicalization (-simplifycfg)
+  // This eases manual annotation of loop bounds as the appearance of
+  // additional nested loops is prohibited by this flag.
+  // When we manage to get rid of loop bound annotations by translating
+  // them to markers and constraints, we can enable this transformation
+  // again.
+  // @see patmos::PatmosBaseTool::ConstructOptJob
+  CmdArgs.push_back("-mllvm");
+  CmdArgs.push_back("-disable-separate-nested-loops");
+
+  // Perform most llvm-opt optimizations at link time (but prepare bitcode with
+  // baseline optimizations)
+  if (Arg *A = Args.getLastArg(options::OPT_O_Group))
+    if (!A->getOption().matches(options::OPT_O0)) {
+      CmdArgs.push_back("-disable-llvm-optzns");
+      CmdArgs.push_back("-enable-llvm-baseline-optzns");
+    }
+}
+
 void Clang::AddWebAssemblyTargetArgs(const ArgList &Args,
                                      ArgStringList &CmdArgs) const {
   // Default to "hidden" visibility.
@@ -2879,6 +2972,9 @@ static bool shouldUseFramePointerForTarget(const ArgList &Args,
     // XCore never wants frame pointers, regardless of OS.
     // WebAssembly never wants frame pointers.
     return false;
+  case llvm::Triple::patmos:
+    // Don't use a frame pointer if optimizing for Patmos, regardless of OS.
+    return !areOptimizationsEnabled(Args);
   default:
     break;
   }
@@ -6520,6 +6616,1262 @@ void hexagon::Linker::ConstructJob(Compilation &C, const JobAction &JA,
 }
 // Hexagon tools end.
 
+static std::string get_patmos_gold(const ToolChain &TC)
+{
+  char *gold_envvar = getenv("PATMOS_GOLD");
+  if (gold_envvar && strcmp(gold_envvar,"")!=0 ) {
+    if (llvm::sys::fs::exists(gold_envvar))
+      return std::string(gold_envvar);
+    else
+      llvm::report_fatal_error("gold linker specified through PATMOS_GOLD "
+                               "environment variable not found.");
+  }
+
+  std::string longname = TC.getTriple().str() + "-ld";
+  std::string tmp( TC.GetProgramPath(longname.c_str()) );
+  if (tmp != longname)
+    return tmp;
+
+  tmp = TC.GetProgramPath("patmos-elf-ld");
+  if (tmp != "patmos-elf-ld")
+    return tmp;
+
+  tmp = TC.GetProgramPath("patmos-gold");
+  if (tmp != "patmos-gold")
+    return tmp;
+
+  tmp = TC.GetProgramPath("patmos-ld");
+  if (tmp != "patmos-ld")
+    return tmp;
+
+  tmp = TC.GetProgramPath("ld.gold");
+  if (tmp != "ld.gold")
+    return tmp;
+
+  tmp = TC.GetProgramPath("ld-new");
+  if (tmp != "ld-new")
+    return tmp;
+
+  return TC.GetProgramPath("ld");
+}
+
+static std::string get_patmos_tool(const ToolChain &TC, StringRef ToolName)
+{
+  std::string longname = (TC.getTriple().str() + "-" + ToolName).str();
+  std::string tmp( TC.GetProgramPath(longname.c_str()) );
+  if (tmp != longname)
+    return tmp;
+
+  std::string PatmosTool = ("patmos-" + ToolName).str();
+  tmp = TC.GetProgramPath(PatmosTool.c_str());
+  if (tmp != PatmosTool)
+    return tmp;
+
+  return TC.GetProgramPath(ToolName.str().c_str());
+}
+
+/// render_patmos_symbol - check if a -mpatmos-<symbol> option was given, if
+/// so render a --defsym to the out arguments list using its value. Otherwise,
+/// render a --defsym using the default value.
+static void render_patmos_symbol(OptSpecifier Opt, const char* Symbol,
+                                 const ArgList &Args, const char *Default,
+                                 ArgStringList &Out)
+{
+  Out.push_back("--defsym");
+  std::string tmp(Symbol);
+  tmp += "=";
+
+  // get option value
+  Arg *a = Args.getLastArg(Opt);
+  tmp += a ? a->getValue() : Default;
+
+  Out.push_back(Args.MakeArgString(tmp));
+}
+
+llvm::sys::fs::file_magic
+patmos::PatmosBaseTool::getFileType(std::string filename) const {
+  llvm::sys::fs::file_magic magic;
+  if (llvm::sys::fs::identify_magic(filename, magic)) {
+    return llvm::sys::fs::file_magic::unknown;
+  }
+
+  return magic;
+}
+
+llvm::sys::fs::file_magic
+patmos::PatmosBaseTool::getBufFileType(const char *buf) const {
+  std::string magic(buf, 4);
+  return llvm::sys::fs::identify_magic(magic);
+}
+
+bool patmos::PatmosBaseTool::isDynamicLibrary(std::string filename) const {
+  llvm::sys::fs::file_magic type = getFileType(filename);
+  switch (type) {
+    default: return false;
+    case llvm::sys::fs::file_magic::macho_fixed_virtual_memory_shared_lib:
+    case llvm::sys::fs::file_magic::macho_dynamically_linked_shared_lib:
+    case llvm::sys::fs::file_magic::macho_dynamically_linked_shared_lib_stub:
+    case llvm::sys::fs::file_magic::elf_shared_object:
+    case llvm::sys::fs::file_magic::pecoff_executable:  return true;
+  }
+}
+
+bool patmos::PatmosBaseTool::isBitcodeArchive(std::string filename) const {
+
+  if (getFileType(filename) != llvm::sys::fs::file_magic::archive) {
+    return false;
+  }
+
+  // check first file in archive if it is a bitcode file
+  auto File = llvm::object::createBinary(filename);
+  if (File.getError()) {
+    return false;
+  }
+
+  if (llvm::object::Archive *a = dyn_cast<llvm::object::Archive>(File.get().getBinary())) {
+    for (llvm::object::Archive::child_iterator i = a->child_begin(), e = a->child_end();
+         i != e; ++i) {
+      // Try opening it as a bitcode file.
+      auto buff = i->get().getMemoryBufferRef();
+      if (buff.getError()) continue;
+      llvm::sys::fs::file_magic FileType = getBufFileType(buff.get().getBufferStart());
+      if (FileType == llvm::sys::fs::file_magic::bitcode) {
+        return true;
+      }
+      if (FileType != llvm::sys::fs::file_magic::unknown) {
+        return false;
+      }
+    }
+  }
+
+  return false;
+}
+
+const char * patmos::PatmosBaseTool::CreateOutputFilename(Compilation &C,
+    const InputInfo &Output, const char * TmpPrefix, const char *Suffix,
+    bool IsLastPass) const
+{
+  const char * filename = NULL;
+
+  const ArgList &Args = C.getArgs();
+  const Driver &D = TC.getDriver();
+
+  if (IsLastPass) {
+    if (Output.isFilename()) {
+      filename = Args.MakeArgString(Output.getFilename());
+    }
+    else {
+      // write to standard-out if nothing is given?!?
+      filename = "-";
+    }
+  } else {
+    if (Args.hasArg(options::OPT_save_temps) && Output.isFilename()) {
+      // take the output's name and append a suffix
+      std::string name(Output.getFilename());
+      filename = Args.MakeArgString((name + "." + Suffix).c_str());
+    }
+    else {
+      StringRef Name = Output.isFilename() ?
+                    llvm::sys::path::filename(Output.getFilename()) : TmpPrefix;
+      std::pair<StringRef, StringRef> Split = Name.split('.');
+      std::string TmpName = D.GetTemporaryPath(Split.first, Suffix);
+      filename = Args.MakeArgString(TmpName.c_str());
+      C.addTempFile(filename);
+    }
+  }
+  return filename;
+}
+
+std::string patmos::PatmosBaseTool::getArgOption(const std::string &Option) const
+{
+  // TODO Check for --<name>=value options??
+
+  if (Option.size() > 2 && Option[2] == '=') {
+    return Option.substr(3);
+  }
+  else if (Option.size() > 2) {
+    return Option.substr(2);
+  }
+  return "";
+}
+
+// Reimplement Linker::FindLib() to search for shared libraries first
+// unless OnlyStatic is true.
+std::string patmos::PatmosBaseTool::FindLib(StringRef LibName,
+                         const std::vector<std::string> &Directories,
+                         bool OnlyStatic) const
+{
+  std::string FilePath(LibName);
+
+  if (llvm::sys::fs::exists(FilePath) &&
+      (isArchive(FilePath) || (!OnlyStatic && isDynamicLibrary(FilePath))))
+    return FilePath;
+
+  // Now iterate over the directories
+  for (std::vector<std::string>::const_iterator Iter = Directories.begin();
+       Iter != Directories.end(); ++Iter) {
+    SmallString<128> FullPath(*Iter);
+
+    llvm::sys::path::append(FullPath, ("lib" + LibName).str());
+
+    // adding a dummy extension so that replace_extension does the right thing
+    FullPath += ".dummy";
+
+    // Either we only want static libraries or we didn't find a
+    // dynamic library so try libX.a
+    llvm::sys::path::replace_extension(FullPath, "a");
+    if (isArchive(FullPath.str()))
+      return FullPath.str();
+
+    // libX.bca
+    llvm::sys::path::replace_extension(FullPath, "bca");
+    if (isArchive(FullPath.str()))
+      return FullPath.str();
+
+    if (!OnlyStatic) {
+      // Try libX.so or libX.dylib
+      // TODO is there a better way to get the shared-lib file extension?
+      llvm::sys::path::replace_extension(FullPath, LTDL_SHLIB_EXT);
+      if (isDynamicLibrary(FullPath.str())) // Native shared library
+        return FullPath.str();
+      if (isBitcodeFile(FullPath.str()))    // .so containing bitcode
+        return FullPath.str();
+    }
+  }
+
+  // No libraries were found
+  return "";
+}
+
+std::vector<std::string>
+patmos::PatmosBaseTool::FindBitcodeLibPaths(const ArgList &Args,
+                                            bool LookupSysPaths) const
+{
+  ArgStringList LibArgs;
+
+  // Use this little trick to prevent duplicating the library path options code
+  AddLibraryPaths(Args, LibArgs, false);
+
+  // To make sure that we catch all -L options of llvm-link, we add -Wl and
+  // -Xlinker.
+  for (ArgList::const_iterator it = Args.begin(), ie = Args.end(); it != ie;
+      ++it)
+  {
+    const Arg *A = *it;
+
+    if (A->getOption().matches(options::OPT_Wl_COMMA) ||
+        A->getOption().matches(options::OPT_Xlinker))
+    {
+      A->renderAsInput(Args, LibArgs);
+    }
+  }
+
+  // Parse all -L options
+  std::vector<std::string> LibPaths;
+  for (ArgStringList::iterator it = LibArgs.begin(), ie = LibArgs.end();
+       it != ie; ++it) {
+    std::string Arg = *it;
+    if (Arg.substr(0, 2) != "-L") continue;
+
+    LibPaths.push_back(getArgOption(Arg));
+  }
+
+  // Collect all the lookup paths
+  if (LookupSysPaths) {
+    // Add same paths as Linker::addSystemPaths().
+    // No use in checking gold linker system paths, if not found we link
+    // using gold in any case.
+    LibPaths.insert(LibPaths.begin(), std::string("./"));
+  }
+
+  return LibPaths;
+}
+
+bool patmos::PatmosBaseTool::isBitcodeOption(StringRef Option,
+                     const std::vector<std::string> &LibPaths) const
+{
+  if (Option.str().substr(0,2) != "-l") {
+    // standard input file
+    return isBitcodeFile(Option);
+  }
+
+  std::string LibName = getArgOption(Option);
+
+  std::string Filename = FindLib(LibName, LibPaths, false);
+  if (!Filename.empty()) {
+    // accept linking with bitcode files
+    return isBitcodeFile(Filename) || isBitcodeArchive(Filename);
+  }
+  return false;
+}
+
+Arg* patmos::PatmosBaseTool::GetOptLevel(const ArgList &Args, char &Lvl) const {
+
+  if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
+    std::string Opt = A->getAsString(Args);
+    if (Opt.length() != 3) {
+      llvm::report_fatal_error("Unsupported optimization option: " + Opt);
+    }
+    Lvl = Opt[2];
+    return A;
+  }
+  return 0;
+}
+
+void patmos::PatmosBaseTool::AddLibraryPaths(const ArgList &Args,
+                                             ArgStringList &CmdArgs,
+                                             bool LinkBinaries) const
+{
+  Args.AddAllArgs(CmdArgs, options::OPT_L);
+
+  addDirectoryList(Args, CmdArgs, "-L", "LIBRARY_PATH");
+
+  // append default search paths
+  const ToolChain::path_list &filePaths = TC.getFilePaths();
+  for(ToolChain::path_list::const_iterator i = filePaths.begin(),
+      ie = filePaths.end(); i != ie; i++) {
+    // construct a library search path
+    std::string path("-L" + *i + "lib/");
+    CmdArgs.push_back(Args.MakeArgString(path.c_str()));
+  }
+}
+
+const char * patmos::PatmosBaseTool::AddInputFiles(const ArgList &Args,
+                           const std::vector<std::string> &LibPaths,
+                           const InputInfoList &Inputs,
+                           ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                           const char *linkedBCFileName,
+                           unsigned &linkedOFileInputPos,
+                           bool AddLibSyms, bool LinkLibraries,
+                           bool HasGoldPass, bool UseLTO) const
+{
+  const Driver &D = TC.getDriver();
+
+  const char *BCOutput = 0;
+
+  for (InputInfoList::const_iterator
+         it = Inputs.begin(), ie = Inputs.end(); it != ie; ++it) {
+    const InputInfo &II = *it;
+
+    //--------------------------------------------------------------------------
+    // handle file inputs
+    if (II.isFilename()) {
+
+      // Check for bitcode files
+      bool IsBitcode = false;
+
+      if (II.getType() == types::TY_AST) {
+        D.Diag(diag::err_drv_no_ast_support) << TC.getTripleString();
+      }
+      else if (II.getType() == types::TY_LLVM_BC ||
+               II.getType() == types::TY_LTO_BC) {
+        IsBitcode = true;
+      }
+      else if (II.getType() == types::TY_Object) {
+
+        llvm::sys::fs::file_magic FT = getFileType(II.getFilename());
+
+        if (FT == llvm::sys::fs::file_magic::archive) {
+          // Should we skip .a files without -l if the do not link in libs?
+          IsBitcode = isBitcodeArchive(II.getFilename());
+        }
+        else if (FT == llvm::sys::fs::file_magic::bitcode) {
+          IsBitcode = true;
+        }
+        else {
+          // Some sort of binary file.. link with gold, even if LTO is not used
+          if (!HasGoldPass) {
+            // TODO use D.Diag() ?
+            llvm::report_fatal_error(Twine(II.getFilename()) +
+                                ": Cannot link binary files when "
+                                "generating bitcode or object file output");
+          }
+        }
+      }
+      else {
+        // Unhandled input file type
+        D.Diag(diag::err_drv_no_linker_llvm_support) << TC.getTripleString();
+      }
+
+      if (IsBitcode && !UseLTO) {
+        if (BCOutput) {
+          // We already have at least one bitcode file, use temp output file
+          BCOutput = linkedBCFileName;
+        } else {
+          // First bitcode file, link compiled file with gold
+          BCOutput = II.getFilename();
+          linkedOFileInputPos = GoldInputs.size();
+        }
+        LinkInputs.push_back(II.getFilename());
+      } else {
+        GoldInputs.push_back(II.getFilename());
+      }
+    }
+    //--------------------------------------------------------------------------
+    // handle -l options
+    else {
+      const Arg &A = II.getInputArg();
+
+      // Reverse translate some rewritten options.
+      if (A.getOption().matches(options::OPT_Z_reserved_lib_stdcxx)) {
+        if (LinkLibraries) {
+          bool IsBitcode = !UseLTO && isBitcodeOption("-lstdc++", LibPaths);
+          if (IsBitcode && !BCOutput) {
+            llvm::report_fatal_error(Twine(II.getAsString()) +
+                                ": Cannot link bitcode library without a "
+                                "previous bitcode input file.");
+          }
+          (IsBitcode ? LinkInputs : GoldInputs).push_back("-lstdc++");
+        }
+        continue;
+      }
+      else if (A.getOption().matches(options::OPT_Wl_COMMA) ||
+               A.getOption().matches(options::OPT_Xlinker) ||
+               A.getOption().matches(options::OPT_Xgold) ||
+               A.getOption().matches(options::OPT_L)) {
+        // already handled
+        continue;
+      }
+      else if (A.getOption().matches(options::OPT_l)) {
+
+        // -l is marked as LinkerInput, so we should always get all -l flags
+        // here, in the correct order.
+        A.claim();
+
+        if (!LinkLibraries) {
+          continue;
+        }
+
+        // -lm is special .. we handle this like a runtime library (should we?)
+        // since we need to link in the libsyms stuff.
+        if (getArgOption(A.getAsString(Args)) == "m") {
+
+          if (AddSystemLibrary(Args, LibPaths, LinkInputs, GoldInputs,
+                              "lib/libmsyms.o", "-lm",
+                              AddLibSyms, HasGoldPass, UseLTO))
+          {
+            if (!BCOutput) {
+              linkedOFileInputPos = GoldInputs.size();
+            }
+            BCOutput = linkedBCFileName;
+          }
+          continue;
+        }
+
+        bool IsBitcode = !UseLTO &&
+                         isBitcodeOption(A.getAsString(Args), LibPaths);
+
+        // Don't render as input
+        A.render(Args, (IsBitcode ? LinkInputs : GoldInputs));
+
+        if (IsBitcode) {
+          if (BCOutput) {
+            // We already have at least one bitcode file, use temp output file
+            BCOutput = linkedBCFileName;
+          } else {
+            // First bitcode file is -l, this will not work
+            llvm::report_fatal_error(Twine(A.getAsString(Args)) +
+                                ": Cannot link bitcode library without a "
+                                "previous bitcode input file.");
+          }
+        }
+      } else {
+        // Uh, what kind of option can this be, and what should we do with it?
+        llvm::report_fatal_error(Twine(A.getAsString(Args)) +
+                            ": unknown linker option.");
+
+      }
+    }
+  }
+
+  return BCOutput;
+}
+
+bool patmos::PatmosBaseTool::AddSystemLibrary(const ArgList &Args,
+                        const std::vector<std::string> &LibPaths,
+                        ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                        const char *libo, const char *libflag,
+                        bool AddLibSyms, bool HasGoldPass, bool UseLTO) const
+{
+  bool IsBitcode = isBitcodeOption(libflag, LibPaths);
+
+  if (IsBitcode && AddLibSyms && libo) {
+    std::string ofile = TC.GetFilePath(libo);
+    (UseLTO ? GoldInputs : LinkInputs).push_back(Args.MakeArgString(ofile));
+  }
+
+  if (libflag) {
+    (UseLTO || !IsBitcode ? GoldInputs : LinkInputs).push_back(libflag);
+  }
+
+  return !UseLTO && IsBitcode && ((AddLibSyms && libo) || libflag);
+}
+
+
+void patmos::PatmosBaseTool::AddStandardLibs(const ArgList &Args,
+                           const std::vector<std::string> &LibPaths,
+                           ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                           bool AddRuntimeLibs, bool AddLibGloss, bool AddLibC,
+                           bool AddLibSyms, StringRef FloatABI,
+                           bool HasGoldPass, bool UseLTO) const
+{
+  // link by default with newlib libc and libpatmos
+  if (AddLibC) {
+    AddSystemLibrary(Args, LibPaths, LinkInputs, GoldInputs,
+                     "lib/libcsyms.o", "-lc",
+                     AddLibSyms, HasGoldPass, UseLTO);
+  }
+
+  // Add support library for newlib libc
+  if (AddLibGloss) {
+    AddSystemLibrary(Args, LibPaths, LinkInputs, GoldInputs,
+                     0, "-lpatmos",
+                     AddLibSyms, HasGoldPass, UseLTO);
+  }
+
+  // link by default with compiler-rt
+  if (AddRuntimeLibs) {
+
+    // softfloat has dependencies to librt, link first
+    if (FloatABI != "hard" && FloatABI != "none") {
+      AddSystemLibrary(Args, LibPaths, LinkInputs, GoldInputs,
+                       "lib/librtsfsyms.o", "-lrtsf",
+                       AddLibSyms, HasGoldPass, UseLTO);
+    }
+
+    AddSystemLibrary(Args, LibPaths, LinkInputs, GoldInputs,
+                     "lib/librtsyms.o", "-lrt",
+                     AddLibSyms, HasGoldPass, UseLTO);
+  }
+
+}
+
+const char * patmos::PatmosBaseTool::PrepareLinkerInputs(const ArgList &Args,
+                     const InputInfoList &Inputs,
+                     ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                     const char *linkedBCFileName,
+                     unsigned &linkedOFileInsertPos,
+                     bool AddStartFiles,
+                     bool AddRuntimeLibs, bool AddLibGloss, bool AddLibC,
+                     bool AddLibSyms, StringRef FloatABI, bool LinkLibraries,
+                     bool HasGoldPass, bool UseLTO) const
+{
+  const char* BCOutput = 0;
+  linkedOFileInsertPos = 0;
+
+  // prepare library lookups
+  // do not add system paths as we call the linker with -nostdlib.
+  std::vector<std::string> BCLibPaths = FindBitcodeLibPaths(Args, false);
+
+  //----------------------------------------------------------------------------
+  // append library search paths (-L) to bitcode linker
+
+  if (LinkLibraries) {
+    AddLibraryPaths(Args, LinkInputs, false);
+  }
+
+  //----------------------------------------------------------------------------
+  // link with start-up files crt0.o and crtbegin.o
+
+  bool AddCrtBeginEnd = AddStartFiles &&
+                        (TC.getTriple().getOS() != llvm::Triple::RTEMS);
+
+  if (AddStartFiles) {
+    std::string Crt0Filename = TC.GetFilePath("lib/crt0.o");
+    std::string CrtBeginFilename = TC.GetFilePath("lib/crtbegin.o");
+
+    if (isBitcodeFile(Crt0Filename)) {
+      const char * crt0 = Args.MakeArgString(Crt0Filename);
+      LinkInputs.push_back(crt0);
+
+      if (AddCrtBeginEnd) {
+        const char * crtbegin = Args.MakeArgString(CrtBeginFilename);
+        LinkInputs.push_back(crtbegin);
+      }
+
+      BCOutput = BCOutput ? linkedBCFileName : crt0;
+    } else {
+      GoldInputs.push_back(Args.MakeArgString(Crt0Filename));
+      linkedOFileInsertPos++;
+
+      if (AddCrtBeginEnd) {
+        GoldInputs.push_back(Args.MakeArgString(CrtBeginFilename));
+        linkedOFileInsertPos++;
+      }
+    }
+  }
+
+  //----------------------------------------------------------------------------
+  // append input files
+
+  const char * InputFile = AddInputFiles(Args, BCLibPaths,
+                Inputs, LinkInputs, GoldInputs,
+                linkedBCFileName, linkedOFileInsertPos,
+                AddLibSyms, LinkLibraries, HasGoldPass, UseLTO);
+
+  if (InputFile) {
+    BCOutput = (BCOutput ? linkedBCFileName : InputFile);
+  }
+
+  //----------------------------------------------------------------------------
+  // link with newlib and compiler-rt libraries
+
+  if (LinkLibraries) {
+    size_t OldSize = LinkInputs.size();
+
+    AddStandardLibs(Args, BCLibPaths, LinkInputs, GoldInputs,
+                    AddRuntimeLibs, AddLibGloss, AddLibC,
+                    AddLibSyms, FloatABI, HasGoldPass, UseLTO);
+
+    // if we added some libs, run link pass
+    if (LinkInputs.size() > OldSize && (BCOutput || AddLibSyms)) {
+      BCOutput = linkedBCFileName;
+    }
+  }
+
+  //----------------------------------------------------------------------------
+  // link with start-up file crtend.o
+
+  if (AddCrtBeginEnd) {
+    std::string CrtEndFilename = TC.GetFilePath("lib/crtend.o");
+
+    if (isBitcodeFile(CrtEndFilename)) {
+      const char * crtend = Args.MakeArgString(CrtEndFilename);
+      LinkInputs.push_back(crtend);
+    } else {
+      GoldInputs.push_back(Args.MakeArgString(CrtEndFilename));
+      linkedOFileInsertPos++;
+    }
+  }
+
+
+  //----------------------------------------------------------------------------
+  // append -L options to gold, but only if we actually need it
+
+  if (LinkLibraries && !GoldInputs.empty()) {
+    ArgStringList TmpArgs;
+
+    AddLibraryPaths(Args, TmpArgs, true);
+
+    GoldInputs.insert(GoldInputs.begin(), TmpArgs.begin(), TmpArgs.end());
+    linkedOFileInsertPos += TmpArgs.size();
+  }
+
+  return BCOutput;
+}
+
+
+void patmos::PatmosBaseTool::ConstructLinkJob(const Tool &Creator,
+                     Compilation &C, const JobAction &JA,
+                     const char *OutputFilename,
+                     const ArgStringList &LinkInputs,
+                     const ArgList &Args) const
+{
+  ArgStringList CmdArgs;
+
+  //----------------------------------------------------------------------------
+  // append linker specific options
+
+  // This must match the argument for FindBitcodeLibPaths in PrepareLinkerInputs
+  CmdArgs.push_back("-nostdlib");
+
+  for (ArgList::const_iterator
+         it = Args.begin(), ie = Args.end(); it != ie; ++it) {
+    Arg *A = *it;
+
+    if (A->getOption().matches(options::OPT_Wl_COMMA) ||
+        A->getOption().matches(options::OPT_Xlinker)) {
+      A->claim();
+      A->renderAsInput(Args, CmdArgs);
+    }
+    else if ((A->getOption().hasFlag(options::LinkerInput) &&
+              !A->getOption().matches(options::OPT_l) &&
+              !A->getOption().matches(options::OPT_Xgold)) ||
+             A->getOption().matches(options::OPT_v))
+    {
+      // It is unfortunate that we have to claim here, as this means
+      // we will basically never report anything interesting for
+      // platforms using a generic gcc, even if we are just using gcc
+      // to get to the assembler.
+      A->claim();
+      A->render(Args, CmdArgs);
+    }
+  }
+
+  //----------------------------------------------------------------------------
+  // append output file for linking
+
+  assert(OutputFilename);
+  CmdArgs.push_back("-o");
+  CmdArgs.push_back(OutputFilename);
+
+  //----------------------------------------------------------------------------
+  // append input arguments
+
+  CmdArgs.append(LinkInputs.begin(), LinkInputs.end());
+
+  //----------------------------------------------------------------------------
+  // execute the linker command
+
+  const char *Exec = Args.MakeArgString(get_patmos_tool(TC, "llvm-link"));
+  C.addCommand(llvm::make_unique<Command>(JA, Creator, Exec, CmdArgs, InputInfo()));
+}
+
+bool patmos::PatmosBaseTool::ConstructOptJob(const Tool &Creator,
+                     Compilation &C, const JobAction &JA,
+                     const char *OutputFilename,
+                     const char *InputFilename,
+                     const ArgList &Args,
+                     bool IsLinkPass, bool LinkAsObject, bool IsLastPass) const
+{
+  ArgStringList OptArgs;
+
+  // TODO add an option (fpatmos-custom-opt-passes=<file>) to specify a file
+  // with a list of passes to enable/disable, or something..
+
+  bool DisableDefaultOpts = Args.hasArg(options::OPT_fpatmos_no_std_link_opts);
+  bool DisableInternalize = LinkAsObject ||
+                   Args.hasArg(options::OPT_fpatmos_disable_internalize);
+
+  int OptLevel = 0;
+  char Lvl;
+  Arg *OptArg;
+  if ((OptArg = GetOptLevel(Args, Lvl))) {
+    switch (Lvl) {
+    case '0': OptLevel = 0; break;
+    case '1': OptLevel = 1; break;
+    case '2': OptLevel = 2; break;
+    case '3': OptLevel = 3; break;
+    // these two need to be > 0, otherwise no opt is triggered
+    case 's': case 'z': OptLevel = 7; break;
+    }
+  }
+
+  //----------------------------------------------------------------------------
+  // append optimization options
+
+  if (IsLinkPass && !DisableDefaultOpts) {
+    if (OptLevel > 0) {
+
+      // adding this here makes the results more similar to plain clang
+      OptArgs.push_back("-sroa");
+
+      // pass -O level to opt verbatim
+      OptArg->renderAsInput(Args, OptArgs);
+
+      if (DisableInternalize) {
+        // works even if in front of -std-link-opts, which adds -internalize
+        OptArgs.push_back("-disable-internalize");
+      } else {
+        // for some reason, we need to add this manually
+        OptArgs.push_back("-internalize");
+        OptArgs.push_back("-globaldce");
+      }
+
+      // @see the note in Clang::AddPatmosTargetArgs()
+      OptArgs.push_back("-disable-separate-nested-loops");
+
+      // simplifying library calls breaks stuff
+      OptArgs.push_back("-disable-simplify-libcalls");
+    }
+
+    if (OptLevel == 3) {
+      // added only for -O3.
+      // (provides -internalize, -globalsmodref-aa over -O3)
+      OptArgs.push_back("-std-link-opts");
+    }
+  }
+
+  // Note: We do not add -O* here, this is done per object by clang.
+
+  // add custom options
+  for (ArgList::const_iterator
+         it = Args.begin(), ie = Args.end(); it != ie; ++it) {
+    Arg *A = *it;
+
+    if (A->getOption().matches(options::OPT_Xopt)) {
+      A->claim();
+      A->renderAsInput(Args, OptArgs);
+    }
+  }
+
+  //----------------------------------------------------------------------------
+  // append output and input files
+
+  if (!IsLastPass && OptArgs.empty()) {
+    return false;
+  }
+
+  assert(OutputFilename);
+  OptArgs.push_back("-o");
+  OptArgs.push_back(OutputFilename);
+
+  OptArgs.push_back(InputFilename);
+
+  //----------------------------------------------------------------------------
+  // execute opt command
+
+  const char *OptExec = Args.MakeArgString(get_patmos_tool(TC, "opt"));
+  C.addCommand(llvm::make_unique<Command>(JA, Creator, OptExec, OptArgs, InputInfo()));
+
+  return true;
+}
+
+void patmos::PatmosBaseTool::ConstructLLCJob(const Tool &Creator,
+    Compilation &C, const JobAction &JA,
+    const char *OutputFilename, const char *InputFilename,
+    const ArgList &Args,
+    bool EmitAsm) const
+{
+  ArgStringList LLCArgs;
+
+  bool ChangedFloatABI;
+  StringRef FloatABI = getPatmosFloatABI(TC.getDriver(), C.getArgs(),
+                                         TC.getTriple(), ChangedFloatABI);
+
+  //----------------------------------------------------------------------------
+  // append -O and -m options
+
+  char OptLevel;
+  if (Arg* A = GetOptLevel(Args, OptLevel)) {
+    switch (OptLevel) {
+    case '0':
+    case '1':
+    case '2':
+    case '3':
+      A->render(Args, LLCArgs);
+      break;
+    default:
+      // LLC does not support -Os, -Oz, ..; uses -O2 instead
+      LLCArgs.push_back("-O2");
+      break;
+    }
+  } else {
+    // If no -O level is supplied, force llc to use -O0
+    LLCArgs.push_back("-O0");
+  }
+
+  // We enable printing labels for all blocks by default in Patmos
+  LLCArgs.push_back("-mforce-block-labels");
+
+  // @see the note in Clang::AddPatmosTargetArgs()
+  LLCArgs.push_back("-disable-separate-nested-loops");
+
+  // floating point arguments are different for LLC
+  for (ArgList::const_iterator
+         it = Args.begin(), ie = Args.end(); it != ie; ++it) {
+    Arg *A = *it;
+
+    if (A->getOption().matches(options::OPT_msoft_float) ||
+        A->getOption().matches(options::OPT_mhard_float) ||
+        A->getOption().matches(options::OPT_mfloat_abi_EQ) ||
+        A->getOption().matches(options::OPT_mno_soft_float) ||
+        A->getOption().matches(options::OPT_m_patmos_Features_Group) ||
+        A->getOption().matches(options::OPT_m_Group))
+    {
+      // Known -m options are already handled in the driver.
+      continue;
+    }
+    else if (A->getOption().matches(options::OPT_m_PML_Group) ||
+             A->getOption().matches(options::OPT_m_patmos_llc_Group))
+    {
+      // PML options and known llc options are directly passed on to LLC.
+      A->claim();
+      A->render(Args, LLCArgs);
+    }
+    else if (A->getSpelling().startswith("-mpatmos"))
+    {
+      // Unknown -mpatmos options are not passed on, make an error out of this!
+      std::string msg = "Use -Xllc " + A->getAsString(Args) +
+                        " to pass this option to llc.";
+      llvm::report_fatal_error(msg);
+    }
+    else if (A->getOption().matches(options::OPT_Xllc)) {
+      A->claim();
+      A->renderAsInput(Args, LLCArgs);
+    }
+  }
+
+  if (ChangedFloatABI) {
+    LLCArgs.push_back("-float-abi");
+    LLCArgs.push_back(FloatABI == "hard" ? "hard" : "soft");
+  }
+
+  //----------------------------------------------------------------------------
+  // generate object file
+
+  // Checking for JA.getType() == types::TY_Image does not tell us if we want to
+  // generate asm code since we told clang that assembly files are .bc files
+  if (EmitAsm) {
+    LLCArgs.push_back("-show-mc-encoding");
+  } else {
+    LLCArgs.push_back("-filetype=obj");
+  }
+
+  //----------------------------------------------------------------------------
+  // append output file for code generation
+
+  assert(OutputFilename);
+  LLCArgs.push_back("-o");
+  LLCArgs.push_back(OutputFilename);
+
+  //----------------------------------------------------------------------------
+  // append linked BC name as input
+
+  LLCArgs.push_back(InputFilename);
+
+  const char *LLCExec = Args.MakeArgString(get_patmos_tool(TC, "llc"));
+  C.addCommand(llvm::make_unique<Command>(JA, Creator, LLCExec, LLCArgs, InputInfo()));
+}
+
+void patmos::PatmosBaseTool::ConstructGoldJob(const Tool &Creator,
+    Compilation &C, const JobAction &JA,
+    const char *OutputFilename, const ArgStringList &GoldInputs,
+    const ArgList &Args, bool UseLTO,
+    bool LinkRelocatable, bool AddStackSymbols) const
+{
+  ArgStringList LDArgs;
+
+  //----------------------------------------------------------------------------
+  // linking options
+
+  if (UseLTO) {
+    // try to add -plugin option, this is actually required
+    std::string plugin( TC.GetProgramPath("LLVMgold.so") );
+    if (plugin != "LLVMgold.so") {
+      LDArgs.push_back("--plugin");
+      LDArgs.push_back(Args.MakeArgString(plugin));
+    }
+  }
+
+  Args.AddAllArgs(LDArgs, options::OPT_T_Group);
+
+  Args.AddAllArgs(LDArgs, options::OPT_e);
+
+  LDArgs.push_back("-nostdlib");
+  LDArgs.push_back("-static");
+
+  if (LinkRelocatable) {
+    // Keep relocations
+    LDArgs.push_back("-r");
+  } else {
+    LDArgs.push_back("--defsym");
+    LDArgs.push_back("__heap_start=end");
+
+    render_patmos_symbol(options::OPT_mpatmos_heap_end,
+                         "__heap_end", Args, "0x100000", LDArgs);
+
+    if (AddStackSymbols) {
+      render_patmos_symbol(options::OPT_mpatmos_shadow_stack_base,
+                           "_shadow_stack_base", Args, "0x1f8000", LDArgs);
+
+      render_patmos_symbol(options::OPT_mpatmos_stack_base,
+                           "_stack_cache_base", Args, "0x200000", LDArgs);
+    }
+  }
+
+  // Do not append arguments given from the Commandline before
+  // setting the defaults
+  for (ArgList::const_iterator
+         it = Args.begin(), ie = Args.end(); it != ie; ++it) {
+    Arg *A = *it;
+
+    if (A->getOption().matches(options::OPT_Xgold)) {
+      A->claim();
+      A->renderAsInput(Args, LDArgs);
+    }
+  }
+
+  if (Args.hasArg(options::OPT_v))
+    LDArgs.push_back("-verbose");
+
+  //----------------------------------------------------------------------------
+  // append output file for code generation
+
+  LDArgs.push_back("-o");
+  LDArgs.push_back(OutputFilename);
+
+  //----------------------------------------------------------------------------
+  // append all linker input arguments and construct the link command
+
+  LDArgs.append(GoldInputs.begin(), GoldInputs.end());
+
+  const char *LDExec = Args.MakeArgString(get_patmos_gold(TC));
+  C.addCommand(llvm::make_unique<Command>(JA, Creator, LDExec, LDArgs, InputInfo()));
+}
+
+
+void patmos::Compile::ConstructJob(Compilation &C, const JobAction &JA,
+                               const InputInfo &Output,
+                               const InputInfoList &Inputs,
+                               const ArgList &Args,
+                               const char *LinkingOutput) const
+{
+  // TODO to make this more standard-compliant, set to false
+  // Note: by default, patmos-clang emits .bc even with no --emit-llvm
+  bool EmitLLVM = true;
+  // only used when EmitLLVM is false
+  bool EmitAsm = false;
+  // do we have a link phase where we call llc
+  bool IsLastPhase = false;
+
+  if (Arg *A = C.getArgs().getLastArg(options::OPT_emit_llvm,
+                                      options::OPT_fpatmos_emit_llvm,
+                                      options::OPT_fpatmos_emit_reloc,
+                                      options::OPT_fpatmos_emit_asm))
+  {
+    if (A->getOption().matches(options::OPT_emit_llvm) ||
+        A->getOption().matches(options::OPT_fpatmos_emit_llvm))
+    {
+      EmitLLVM = true;
+    }
+    else if (A->getOption().matches(options::OPT_fpatmos_emit_reloc)) {
+      EmitLLVM = false;
+      EmitAsm = false;
+    }
+    else if (A->getOption().matches(options::OPT_fpatmos_emit_asm)) {
+      EmitLLVM = false;
+      EmitAsm = true;
+    }
+    A->claim();
+  }
+  if (C.getArgs().hasArg(options::OPT_S))
+  {
+    EmitAsm = true;
+    IsLastPhase = true;
+  }
+  if (C.getArgs().hasArg(options::OPT_c))
+  {
+    EmitAsm = false;
+    IsLastPhase = true;
+  }
+
+  // TODO instead of running LLC separately, we might also just add the opt/LLC
+  //      options to the clang options using -mllvm for -Xopt/-Xllc options and
+  //      let clang -cc1 do the work (but this is slightly harder to debug).
+
+  // If this is not the last phase or if we emit llvm-code, we just call clang
+  if (EmitLLVM || !IsLastPhase) {
+
+    Clang::ConstructJob(C, JA, Output, Inputs, Args, LinkingOutput);
+
+  } else {
+    // Run llc afterwards, no linking phase
+    const char *BCFilename = CreateOutputFilename(C, Output, "clang-",
+                                                  "bc", false);
+
+    const InputInfo TmpOutput(types::TY_LLVM_BC, BCFilename,
+                              Inputs[0].getFilename());
+
+    //--------------------------------------------------------------------------
+    // Run clang
+
+    Clang::ConstructJob(C, JA, TmpOutput, Inputs, Args, LinkingOutput);
+
+    //--------------------------------------------------------------------------
+    // Run LLC on output
+
+    ConstructLLCJob(*this, C, JA, Output.getFilename(), BCFilename, Args,
+                    EmitAsm);
+
+  }
+}
+
+
+void patmos::Link::ConstructJob(Compilation &C, const JobAction &JA,
+                               const InputInfo &Output,
+                               const InputInfoList &Inputs,
+                               const ArgList &Args,
+                               const char *LinkingOutput) const
+{
+  const ToolChain &TC = getToolChain();
+  const Driver &D = TC.getDriver();
+
+  // In the link phase, we do the following:
+  // - run llvm-link, link in bitcode files:
+  //    - startup files if not -nostartfiles
+  //    - all bitcode input files
+  //    - all -l<library>, except when linking as object file
+  //    - all standard libraries, except when disabled
+  //    - all runtime libs, except when disabled
+  //    - all required libsyms.o files for linked in libs, except when disabled
+  // - run llc on result, perform optimization on linked bitcode
+  // - run gold on result, link in:
+  //    - all ELF input files, same order as llvm-link
+
+  //----------------------------------------------------------------------------
+  // read out various command line options
+
+  bool LinkRTEMS = (TC.getTriple().getOS() == llvm::Triple::RTEMS);
+
+  bool ChangedFloatABI;
+  StringRef FloatABI = getPatmosFloatABI(TC.getDriver(), C.getArgs(),
+                                         TC.getTriple(), ChangedFloatABI);
+
+  // add lib*syms.o options to llvm-link
+  bool AddLibSyms = !C.getArgs().hasArg(options::OPT_nolibsyms);
+  // add crt0
+  bool AddStartFiles = !C.getArgs().hasArg(options::OPT_nostartfiles);
+  // add librt, librtsf
+  bool AddRuntimeLibs = !C.getArgs().hasArg(options::OPT_noruntimelibs);
+  // add libpatmos, libc, ..
+  bool AddStdLibs = !C.getArgs().hasArg(options::OPT_nostdlib);
+  // add libc
+  bool AddLibC = !C.getArgs().hasArg(options::OPT_nolibc) && AddStdLibs;
+  // add support library for newlib libc
+  bool AddLibGloss = AddStdLibs && !LinkRTEMS;
+  // add any default libs at all?
+  bool AddDefaultLibs = !C.getArgs().hasArg(options::OPT_nodefaultlibs);
+
+  // Do not link in -l, do not link in startup code or standard-libs
+  bool LinkAsObject = C.getArgs().hasArg(options::OPT_fpatmos_link_object);
+
+  // Do not execute llc and gold
+  bool EmitLLVM = C.getArgs().hasArg(options::OPT_fpatmos_emit_llvm) ||
+                  C.getArgs().hasArg(options::OPT_emit_llvm);
+  // Do not execute gold
+  bool EmitObject = C.getArgs().hasArg(options::OPT_fpatmos_emit_reloc);
+  // Do not execute gold, emit assembler
+  bool EmitAsm = C.getArgs().hasArg(options::OPT_fpatmos_emit_asm);
+
+  // do not run opt after bitcode linking
+  bool SkipOpt = C.getArgs().hasArg(options::OPT_fpatmos_skip_opt);
+
+  // link all -l and ELF .o files with gold and libLTO plugin
+  bool UseLTO = C.getArgs().hasArg(options::OPT_flto);
+
+  if (!AddDefaultLibs) {
+    AddRuntimeLibs = false;
+    AddStdLibs = false;
+    AddLibC = false;
+    AddLibGloss = false;
+  }
+  if (LinkAsObject) {
+    AddStartFiles = false;
+    AddLibSyms = false;
+    AddRuntimeLibs = false;
+    AddStdLibs = false;
+    AddLibC = false;
+    AddLibGloss = false;
+    EmitLLVM = !(EmitObject || EmitAsm);
+  }
+
+  bool StopAfterLink = EmitLLVM && SkipOpt;
+  bool StopAfterOpt = EmitLLVM;
+  bool StopAfterLLC = EmitAsm;
+
+  //----------------------------------------------------------------------------
+  // Sanity checks and check for some unsupported options
+
+  if (EmitObject && EmitAsm) {
+    llvm::report_fatal_error("-fpatmos-emit-reloc and -fpatmos-emit-asm are mutually exclusive");
+  }
+  if (EmitLLVM && (EmitObject || EmitAsm)) {
+    llvm::report_fatal_error("-fpatmos-emit-llvm cannot be used with -fpatmos-emit-reloc or -fpatmos-emit-asm");
+  }
+  if (UseLTO && LinkAsObject) {
+    llvm::report_fatal_error("-fpatmos-link-object and linking with LTO support is mutually exclusive.");
+  }
+  if (UseLTO && (EmitLLVM || EmitObject || EmitAsm)) {
+    llvm::report_fatal_error("-fpatmos-emit-* cannot be used when linking with LTO support");
+  }
+
+  if (Arg *A = Args.getLastArg(options::OPT_shared)) {
+    D.Diag(diag::err_drv_unsupported_opt) << A->getAsString(Args);
+  }
+  if (Arg *A = Args.getLastArg(options::OPT_static)) {
+    D.Diag(diag::err_drv_unsupported_opt) << A->getAsString(Args);
+  }
+
+  //////////////////////////////////////////////////////////////////////////////
+  // Prepare linker arguments
+
+  const char *linkedBCFileName =
+           CreateOutputFilename(C, Output, "link-", "bc", StopAfterLink);
+
+  ArgStringList LinkInputs, GoldInputs;
+  unsigned linkedOFileInsertPos;
+
+  bool HasGoldPass = !StopAfterLink && !EmitLLVM && !StopAfterLLC;
+
+  const char *BCFile = PrepareLinkerInputs(Args, Inputs,
+                         LinkInputs, GoldInputs,
+                         linkedBCFileName, linkedOFileInsertPos,
+                         AddStartFiles, AddRuntimeLibs, AddLibGloss, AddLibC,
+                         AddLibSyms, FloatABI,
+                         !LinkAsObject, HasGoldPass, UseLTO);
+
+  bool RequiresLink = true;
+  if ((!BCFile || BCFile != linkedBCFileName) && !StopAfterLink) {
+    // No bitcode input, or only a single bitcode input file, skip link
+    // pass and use input directly if it is a single file.
+    RequiresLink = false;
+    linkedBCFileName = BCFile;
+  }
+
+  //////////////////////////////////////////////////////////////////////////////
+  // build LINK command
+
+  if (RequiresLink) {
+    ConstructLinkJob(*this, C, JA, linkedBCFileName, LinkInputs, Args);
+  }
+
+  if (StopAfterLink) {
+    return;
+  }
+
+  //////////////////////////////////////////////////////////////////////////////
+  // build OPT command
+
+  if (!SkipOpt && linkedBCFileName) {
+    char const *optimizedBCFileName =
+                CreateOutputFilename(C, Output, "opt-", "opt.bc", StopAfterOpt);
+
+    if (ConstructOptJob(*this, C, JA, optimizedBCFileName, linkedBCFileName,
+                        Args, true, LinkAsObject, StopAfterOpt)) {
+      linkedBCFileName = optimizedBCFileName;
+    }
+  }
+
+  // If we only want to emit bitcode, we are done now.
+  if (StopAfterOpt) {
+    return;
+  }
+
+  //////////////////////////////////////////////////////////////////////////////
+  // build LLC command
+
+  const char *linkedOFileName = EmitLLVM ? 0 :
+           CreateOutputFilename(C, Output, "llc-", "bc.o", StopAfterLLC);
+
+  if (linkedBCFileName) {
+    ConstructLLCJob(*this, C, JA, linkedOFileName, linkedBCFileName, Args,
+                    EmitAsm);
+  }
+
+  // If we do not want to create an executable file, we are done now
+  if (StopAfterLLC) {
+    return;
+  }
+
+  //////////////////////////////////////////////////////////////////////////////
+  // build LD command
+
+  char const *linkedELFFileName = CreateOutputFilename(C, Output, "gold-",
+                                                       ".out", true);
+
+  if (linkedBCFileName) {
+    // If we compiled a bitcode file, insert the compiled file into gold args
+    GoldInputs.insert(GoldInputs.begin() + linkedOFileInsertPos,
+                      linkedOFileName);
+  }
+
+  ConstructGoldJob(*this, C, JA, linkedELFFileName, GoldInputs, Args,
+                   UseLTO, EmitObject || LinkAsObject, !LinkRTEMS);
+}
+// Patmos tools end.
+
+
 void amdgpu::Linker::ConstructJob(Compilation &C, const JobAction &JA,
                                   const InputInfo &Output,
                                   const InputInfoList &Inputs,
diff --git a/lib/Driver/Tools.h b/lib/Driver/Tools.h
index 7ed2cfd091..e775ceaad4 100644
--- a/lib/Driver/Tools.h
+++ b/lib/Driver/Tools.h
@@ -84,6 +84,8 @@ private:
                         llvm::opt::ArgStringList &CmdArgs) const;
   void AddHexagonTargetArgs(const llvm::opt::ArgList &Args,
                             llvm::opt::ArgStringList &CmdArgs) const;
+  void AddPatmosTargetArgs(const llvm::opt::ArgList &Args,
+                           llvm::opt::ArgStringList &CmdArgs) const;
   void AddWebAssemblyTargetArgs(const llvm::opt::ArgList &Args,
                                 llvm::opt::ArgStringList &CmdArgs) const;
 
@@ -238,6 +240,165 @@ public:
 };
 } // end namespace hexagon.
 
+namespace patmos {
+
+  class PatmosBaseTool {
+    const ToolChain &TC;
+    std::vector<std::string> LibraryPaths;
+  public:
+    PatmosBaseTool(const ToolChain &TC) : TC(TC) {}
+
+  protected:
+    // Some helper methods to construct arguments in ConstructJob
+    llvm::sys::fs::file_magic getFileType(std::string filename) const;
+
+    llvm::sys::fs::file_magic getBufFileType(const char *buf) const;
+
+    bool isBitcodeFile(std::string filename) const {
+      return getFileType(filename) == llvm::sys::fs::file_magic::bitcode;
+    }
+
+    bool isArchive(std::string filename) const {
+      return getFileType(filename) == llvm::sys::fs::file_magic::archive;
+    }
+
+    bool isDynamicLibrary(std::string filename) const;
+
+    bool isBitcodeArchive(std::string filename) const;
+
+    bool isBitcodeOption(StringRef Option,
+                         const std::vector<std::string> &LibPaths) const;
+
+    const char * CreateOutputFilename(Compilation &C, const InputInfo &Output,
+                                      const char * TmpPrefix,
+                                      const char *Suffix,
+                                      bool IsLastPass) const;
+
+    /// Get the option value of an argument
+    std::string getArgOption(const std::string &Arg) const;
+
+    std::string FindLib(StringRef LibName,
+                        const std::vector<std::string> &Directories,
+                        bool OnlyStatic) const;
+
+    std::vector<std::string> FindBitcodeLibPaths(const llvm::opt::ArgList &Args,
+                                                 bool LookupSysPaths) const;
+
+    /// Get the last -O<Lvl> optimization level specifier. If no -O option is
+    /// given, return NULL.
+    llvm::opt::Arg* GetOptLevel(const llvm::opt::ArgList &Args, char &Lvl) const;
+
+    /// Add -L arguments
+    void AddLibraryPaths(const llvm::opt::ArgList &Args, ArgStringList &CmdArgs,
+                         bool LinkBinaries) const;
+
+    /// The HasGoldPass arguments tells the function if
+    /// we will execute gold or if linking with ELFs should throw an error.
+    /// Return the
+    const char * AddInputFiles(const llvm::opt::ArgList &Args,
+                       const std::vector<std::string> &LibPaths,
+                       const InputInfoList &Inputs,
+                       ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                       const char *linkedBCFileName,
+                       unsigned &linkedOFileInputPos,
+                       bool AddLibSyms, bool LinkLibraries,
+                       bool HasGoldPass, bool UseLTO) const;
+
+    /// Return true if any options have been added to LinkInputs.
+    bool AddSystemLibrary(const llvm::opt::ArgList &Args,
+                          const std::vector<std::string> &LibPaths,
+                          ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                          const char *libo, const char *libflag,
+                          bool AddLibSyms, bool HasGoldPass, bool UseLTO) const;
+
+    /// Add arguments to link with libc, librt, librtsf, libpatmos
+    void AddStandardLibs(const llvm::opt::ArgList &Args,
+                         const std::vector<std::string> &LibPaths,
+                         ArgStringList &LinkInputs, ArgStringList &GoldInputs,
+                         bool AddRuntimeLibs, bool AddLibGloss, bool AddLibC,
+                         bool AddLibSyms, StringRef FloatABI,
+                         bool HasGoldPass, bool UseLTO) const;
+
+
+    /// Returns linkedBCFileName if files need to be linked, or the filename of
+    /// the only bitcode input file if there is no need to link, or null if
+    /// there are no bitcode inputs.
+    /// @linkedOFileInsertPos - position in GoldInputs where to insert the
+    /// compiled bitcode file into.
+    const char * PrepareLinkerInputs(const llvm::opt::ArgList &Args,
+                         const InputInfoList &Inputs,
+                         llvm::opt::ArgStringList &LinkInputs,
+                         llvm::opt::ArgStringList &GoldInputs,
+                         const char *linkedBCFileName,
+                         unsigned &linkedOFileInsertPos,
+                         bool AddStartFiles,
+                         bool AddRuntimeLibs, bool AddLibGloss, bool AddLibC,
+                         bool AddLibSyms, StringRef FloatABI,
+                         bool LinkLibraries,
+                         bool HasGoldPass, bool UseLTO) const;
+
+    void ConstructLinkJob(const Tool &Creator, Compilation &C,
+                          const JobAction &JA,
+                          const char *OutputFilename,
+                          const llvm::opt::ArgStringList &LinkInputs,
+                          const llvm::opt::ArgList &TCArgs) const;
+
+    // Construct an optimization job
+    // @IsLinkPass - If true, add standard link optimizations
+    bool ConstructOptJob(const Tool &Creator, Compilation &C,
+                         const JobAction &JA,
+                         const char *OutputFilename,
+                         const char *InputFilename,
+                         const llvm::opt::ArgList &TCArgs,
+                         bool IsLinkPass, bool LinkAsObject, bool IsLastPass) const;
+
+    void ConstructLLCJob(const Tool &Creator, Compilation &C,
+                      const JobAction &JA,
+                      const char *OutputFilename,
+                      const char *InputFilename,
+                      const llvm::opt::ArgList &TCArgs,
+                      bool EmitAsm) const;
+
+    void ConstructGoldJob(const Tool &Creator, Compilation &C,
+                          const JobAction &JA,
+                          const char *OutputFilename,
+                          const llvm::opt::ArgStringList &GoldInputs,
+                          const llvm::opt::ArgList &TCArgs, bool UseLTO,
+                          bool LinkRelocatable, bool AddStackSymbols) const;
+  };
+
+  class LLVM_LIBRARY_VISIBILITY Compile : public Clang, protected PatmosBaseTool
+  {
+  public:
+    Compile(const ToolChain &TC) : Clang(TC), PatmosBaseTool(TC) {}
+
+    void ConstructJob(Compilation &C, const JobAction &JA,
+                      const InputInfo &Output,
+                      const InputInfoList &Inputs,
+                      const llvm::opt::ArgList &TCArgs,
+                      const char *LinkingOutput) const override;
+
+    bool hasIntegratedAssembler() const override { return false; }
+  };
+
+  class LLVM_LIBRARY_VISIBILITY Link : public Tool, protected PatmosBaseTool {
+  public:
+    Link(const ToolChain &TC) : Tool("patmos::Link",
+                                     "link  via llvm-link, opt and llc", TC),
+                                PatmosBaseTool(TC) {}
+
+    void ConstructJob(Compilation &C, const JobAction &JA,
+                      const InputInfo &Output,
+                      const InputInfoList &Inputs,
+                      const llvm::opt::ArgList &TCArgs,
+                      const char *LinkingOutput) const override;
+
+    bool hasIntegratedCPP() const override { return false; }
+    bool isLinkJob() const override { return true; }
+
+  };
+} // end namespace patmos
+
 namespace amdgpu {
 
 class LLVM_LIBRARY_VISIBILITY Linker : public GnuTool {
diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index 31da8bc721..3715737e95 100644
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -1061,6 +1061,7 @@ static InputKind ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args,
   Opts.ASTDumpLookups = Args.hasArg(OPT_ast_dump_lookups);
   Opts.UseGlobalModuleIndex = !Args.hasArg(OPT_fno_modules_global_index);
   Opts.GenerateGlobalModuleIndex = Opts.UseGlobalModuleIndex;
+  Opts.FlowfactExportFile = Args.getLastArgValue(OPT_flowfact_export_EQ);
   Opts.ModuleMapFiles = Args.getAllArgValues(OPT_fmodule_map_file);
   Opts.ModuleFiles = Args.getAllArgValues(OPT_fmodule_file);
   Opts.ModulesEmbedFiles = Args.getAllArgValues(OPT_fmodules_embed_file_EQ);
diff --git a/lib/Frontend/FrontendActions.cpp b/lib/Frontend/FrontendActions.cpp
index 407ccea2e7..7236cc3ed7 100644
--- a/lib/Frontend/FrontendActions.cpp
+++ b/lib/Frontend/FrontendActions.cpp
@@ -45,7 +45,6 @@ void InitOnlyAction::ExecuteAction() {
 //===----------------------------------------------------------------------===//
 // AST Consumer Actions
 //===----------------------------------------------------------------------===//
-
 std::unique_ptr<ASTConsumer>
 ASTPrintAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) {
   if (raw_ostream *OS = CI.createDefaultOutputFile(false, InFile))
@@ -76,6 +75,12 @@ DeclContextPrintAction::CreateASTConsumer(CompilerInstance &CI,
   return CreateDeclContextPrinter();
 }
 
+std::unique_ptr<ASTConsumer>
+FlowfactExportAction::CreateASTConsumer(CompilerInstance &CI,
+                                        StringRef InFile) {
+  return CreateFlowfactExporter(CI.getFrontendOpts().FlowfactExportFile);
+}
+
 std::unique_ptr<ASTConsumer>
 GeneratePCHAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) {
   std::string Sysroot;
diff --git a/lib/Parse/ParsePragma.cpp b/lib/Parse/ParsePragma.cpp
index bc70942851..85db7610fe 100644
--- a/lib/Parse/ParsePragma.cpp
+++ b/lib/Parse/ParsePragma.cpp
@@ -17,6 +17,7 @@
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Parse/ParseDiagnostic.h"
 #include "clang/Parse/Parser.h"
+#include "clang/Sema/Loopbound.h"
 #include "clang/Sema/LoopHint.h"
 #include "clang/Sema/Scope.h"
 #include "llvm/ADT/StringSwitch.h"
@@ -144,6 +145,18 @@ private:
   Sema &Actions;
 };
 
+struct PragmaLoopboundHandler : public PragmaHandler {
+  PragmaLoopboundHandler() : PragmaHandler("loopbound") {}
+  void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+                    Token &FirstToken) override;
+};
+
+struct PragmaPlatinHandler : public PragmaHandler {
+  PragmaPlatinHandler() : PragmaHandler("platin") { }
+  void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+                    Token &FirstToken) override;
+};
+
 struct PragmaLoopHintHandler : public PragmaHandler {
   PragmaLoopHintHandler() : PragmaHandler("loop") {}
   void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
@@ -233,6 +246,12 @@ void Parser::initializePragmaHandlers() {
   OptimizeHandler.reset(new PragmaOptimizeHandler(Actions));
   PP.AddPragmaHandler("clang", OptimizeHandler.get());
 
+  LoopboundHandler.reset(new PragmaLoopboundHandler());
+  PP.AddPragmaHandler(LoopboundHandler.get());
+
+  PlatinHandler.reset(new PragmaPlatinHandler());
+  PP.AddPragmaHandler(PlatinHandler.get());
+
   LoopHintHandler.reset(new PragmaLoopHintHandler());
   PP.AddPragmaHandler("clang", LoopHintHandler.get());
 
@@ -304,6 +323,12 @@ void Parser::resetPragmaHandlers() {
   PP.RemovePragmaHandler("clang", OptimizeHandler.get());
   OptimizeHandler.reset();
 
+  PP.RemovePragmaHandler(PlatinHandler.get());
+  PlatinHandler.reset();
+
+  PP.RemovePragmaHandler(LoopboundHandler.get());
+  LoopboundHandler.reset();
+
   PP.RemovePragmaHandler("clang", LoopHintHandler.get());
   LoopHintHandler.reset();
 
@@ -760,6 +785,37 @@ bool Parser::HandlePragmaMSInitSeg(StringRef PragmaName,
   return true;
 }
 
+namespace {
+struct PragmaLoopboundInfo {
+  Token PragmaName;
+  Token Min;
+  Token Max;
+};
+} // end anonymous namespace
+
+void Parser::HandlePragmaLoopbound(Loopbound &LB) {
+  assert(Tok.is(tok::annot_pragma_loopbound));
+
+  PragmaLoopboundInfo *Info =
+      static_cast<PragmaLoopboundInfo *>(Tok.getAnnotationValue());
+  ConsumeToken(); // The annotation token.
+
+  LB.PragmaNameLoc = IdentifierLoc::create(
+      Actions.Context,
+      Info->PragmaName.getLocation(),
+      Info->PragmaName.getIdentifierInfo()
+      );
+
+  assert(Info->Min.is(tok::numeric_constant));
+  LB.MinExpr = Actions.ActOnNumericConstant(Info->Min).get();
+
+  assert(Info->Max.is(tok::numeric_constant));
+  LB.MaxExpr = Actions.ActOnNumericConstant(Info->Max).get();
+
+  LB.Range =
+        SourceRange(Info->PragmaName.getLocation(), Info->Max.getLocation());
+}
+
 namespace {
 struct PragmaLoopHintInfo {
   Token PragmaName;
@@ -1910,6 +1966,90 @@ void PragmaOptimizeHandler::HandlePragma(Preprocessor &PP,
   Actions.ActOnPragmaOptimize(IsOn, FirstToken.getLocation());
 }
 
+// #pragma loopbound min NUM max NUM
+void PragmaLoopboundHandler::HandlePragma(Preprocessor &PP,
+                                          PragmaIntroducerKind Introducer,
+                                          Token &LoopboundTok) {
+  Token Tok;
+
+  PragmaLoopboundInfo *Info =
+    new (PP.getPreprocessorAllocator()) PragmaLoopboundInfo;
+  Info->PragmaName = LoopboundTok;
+
+  PP.LexUnexpandedToken(Tok);
+  if (Tok.isNot(tok::identifier) ||
+      !Tok.getIdentifierInfo()->isStr("min")) {
+    PP.Diag(Tok.getLocation(), diag::err_pragma_loopbound_malformed);
+    return;
+  }
+
+  PP.Lex(Tok); // allow macro expansion for minimum
+  if (!Tok.is(tok::numeric_constant)) {
+    PP.Diag(Tok.getLocation(), diag::err_pragma_loopbound_malformed);
+    return;
+  }
+  // store loopbound min
+  Info->Min = Tok;
+
+  PP.LexUnexpandedToken(Tok);
+  if (Tok.isNot(tok::identifier) ||
+      !Tok.getIdentifierInfo()->isStr("max")) {
+    PP.Diag(Tok.getLocation(), diag::err_pragma_loopbound_malformed);
+    return;
+  }
+
+  PP.Lex(Tok); // allow macro expansion for maximum
+  if (!Tok.is(tok::numeric_constant)) {
+    PP.Diag(Tok.getLocation(), diag::err_pragma_loopbound_malformed);
+    return;
+  }
+  // store loopbound max
+  Info->Max = Tok;
+
+  // eat the max
+  PP.Lex(Tok);
+  if (Tok.isNot(tok::eod)) {
+    PP.Diag(Tok.getLocation(), diag::err_pragma_loopbound_malformed);
+    return;
+  }
+
+  // Generate the hint token.
+  Token *TokenArray = new Token[1];
+  TokenArray[0].startToken();
+  TokenArray[0].setKind(tok::annot_pragma_loopbound);
+  TokenArray[0].setLocation(LoopboundTok.getLocation());
+  TokenArray[0].setAnnotationValue(static_cast<void *>(Info));
+  PP.EnterTokenStream(TokenArray, 1, /*DisableMacroExpansion=*/false,
+                      /*OwnsTokens=*/true);
+}
+
+void
+PragmaPlatinHandler::HandlePragma(Preprocessor &PP,
+                                  PragmaIntroducerKind Introducer,
+                                  Token &FirstTok) {
+  SmallVector<Token, 16> Pragma;
+  Token Tok;
+  Tok.startToken();
+  Tok.setKind(tok::annot_pragma_platinff);
+  Tok.setLocation(FirstTok.getLocation());
+
+  while (Tok.isNot(tok::eod)) {
+    Pragma.push_back(Tok);
+    PP.Lex(Tok);
+  }
+  SourceLocation EodLoc = Tok.getLocation();
+  Tok.startToken();
+  Tok.setKind(tok::annot_pragma_platinff_end);
+  Tok.setLocation(EodLoc);
+  Pragma.push_back(Tok);
+
+  Token *Toks = new Token[Pragma.size()];
+  std::copy(Pragma.begin(), Pragma.end(), Toks);
+  PP.EnterTokenStream(Toks, Pragma.size(),
+                      /*DisableMacroExpansion=*/true, /*OwnsTokens=*/true);
+}
+
+
 /// \brief Parses loop or unroll pragma hint value and fills in Info.
 static bool ParseLoopHintValue(Preprocessor &PP, Token &Tok, Token PragmaName,
                                Token Option, bool ValueInParens,
diff --git a/lib/Parse/ParseStmt.cpp b/lib/Parse/ParseStmt.cpp
index edf0dda7df..3993b51573 100644
--- a/lib/Parse/ParseStmt.cpp
+++ b/lib/Parse/ParseStmt.cpp
@@ -369,6 +369,14 @@ Retry:
     HandlePragmaMSVtorDisp();
     return StmtEmpty();
 
+  case tok::annot_pragma_loopbound:
+    ProhibitAttributes(Attrs);
+    return ParsePragmaLoopbound(Stmts, Allowed, TrailingElseLoc, Attrs);
+
+  case tok::annot_pragma_platinff:
+    ProhibitAttributes(Attrs);
+    return ParsePlatinPragma();
+
   case tok::annot_pragma_loop_hint:
     ProhibitAttributes(Attrs);
     return ParsePragmaLoopHint(Stmts, Allowed, TrailingElseLoc, Attrs);
@@ -1878,6 +1886,34 @@ StmtResult Parser::ParseReturnStatement() {
   return Actions.ActOnReturnStmt(ReturnLoc, R.get(), getCurScope());
 }
 
+StmtResult Parser::ParsePragmaLoopbound(StmtVector &Stmts,
+                                        AllowedContsructsKind Allowed,
+                                        SourceLocation *TrailingElseLoc,
+                                        ParsedAttributesWithRange &Attrs) {
+  // Create temporary attribute list.
+  ParsedAttributesWithRange TempAttrs(AttrFactory);
+
+  // Get loopbound and consume annotated token.
+  while (Tok.is(tok::annot_pragma_loopbound)) {
+    Loopbound LB;
+    HandlePragmaLoopbound(LB);
+
+    ArgsUnion ArgLB[] = {ArgsUnion(LB.MinExpr), ArgsUnion(LB.MaxExpr)};
+    TempAttrs.addNew(LB.PragmaNameLoc->Ident, LB.Range, NULL,
+                     LB.PragmaNameLoc->Loc, ArgLB, 2,
+                     AttributeList::AS_Pragma);
+  }
+
+  // Get the next statement.
+  MaybeParseCXX11Attributes(Attrs);
+
+  StmtResult S = ParseStatementOrDeclarationAfterAttributes(
+      Stmts, Allowed, TrailingElseLoc, Attrs);
+
+  Attrs.takeAllFrom(TempAttrs);
+  return S;
+}
+
 StmtResult Parser::ParsePragmaLoopHint(StmtVector &Stmts,
                                        AllowedContsructsKind Allowed,
                                        SourceLocation *TrailingElseLoc,
diff --git a/lib/Parse/Parser.cpp b/lib/Parse/Parser.cpp
index defb74b009..ccefb3dd3f 100644
--- a/lib/Parse/Parser.cpp
+++ b/lib/Parse/Parser.cpp
@@ -85,9 +85,6 @@ Parser::Parser(Preprocessor &pp, Sema &actions, bool skipFunctionBodies)
   // destructor.
   initializePragmaHandlers();
 
-  LoopboundHandler.reset(new PragmaLoopboundHandler());
-  PP.AddPragmaHandler(LoopboundHandler.get());
-
   CommentSemaHandler.reset(new ActionCommentHandler(actions));
   PP.addCommentHandler(CommentSemaHandler.get());
 
@@ -434,9 +431,6 @@ Parser::~Parser() {
 
   resetPragmaHandlers();
 
-  PP.RemovePragmaHandler(LoopboundHandler.get());
-  LoopboundHandler.reset();
-
   PP.removeCommentHandler(CommentSemaHandler.get());
 
   PP.clearCodeCompletionHandler();
diff --git a/lib/Sema/SemaStmtAttr.cpp b/lib/Sema/SemaStmtAttr.cpp
index 984bd078fa..52e19fa354 100644
--- a/lib/Sema/SemaStmtAttr.cpp
+++ b/lib/Sema/SemaStmtAttr.cpp
@@ -43,6 +43,39 @@ static Attr *handleFallThroughAttr(Sema &S, Stmt *St, const AttributeList &A,
                                            A.getAttributeSpellingListIndex());
 }
 
+static Attr *handleLoopboundAttr(Sema &S, Stmt *St, const AttributeList &A,
+                                 SourceRange Range) {
+  Expr *MinExpr = A.getArgAsExpr(0);
+  Expr *MaxExpr = A.getArgAsExpr(1);
+  if (St->getStmtClass() != Stmt::DoStmtClass &&
+      St->getStmtClass() != Stmt::ForStmtClass &&
+      St->getStmtClass() != Stmt::CXXForRangeStmtClass &&
+      St->getStmtClass() != Stmt::WhileStmtClass) {
+    S.Diag(St->getLocStart(), diag::err_pragma_loop_precedes_nonloop)
+       << "#pragma loopbound";
+    return 0;
+  }
+
+  llvm::APSInt MinAPS, MaxAPS;
+  assert(MinExpr != NULL && MaxExpr != NULL);
+  if (!MinExpr->isIntegerConstantExpr(MinAPS, S.Context) ||
+      !MaxExpr->isIntegerConstantExpr(MaxAPS, S.Context)) {
+    S.Diag(A.getLoc(), diag::err_pragma_loopbound_invalid_values);
+    return 0;
+  }
+
+  int MinInt = MinAPS.getSExtValue();
+  int MaxInt = MaxAPS.getSExtValue();
+
+  if ( MinInt < 0 || MaxInt < 0 || MinInt > MaxInt) {
+    S.Diag(A.getLoc(), diag::err_pragma_loopbound_invalid_values);
+    return 0;
+  }
+
+  return ::new (S.Context) LoopboundAttr(A.getRange(), S.Context,
+      MinInt, MaxInt, A.getAttributeSpellingListIndex());
+}
+
 static Attr *handleLoopHintAttr(Sema &S, Stmt *St, const AttributeList &A,
                                 SourceRange) {
   IdentifierLoc *PragmaNameLoc = A.getArgAsIdent(0);
@@ -213,6 +246,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const AttributeList &A,
     return nullptr;
   case AttributeList::AT_FallThrough:
     return handleFallThroughAttr(S, St, A, Range);
+  case AttributeList::AT_Loopbound:
+    return handleLoopboundAttr(S, St, A, Range);
   case AttributeList::AT_LoopHint:
     return handleLoopHintAttr(S, St, A, Range);
   default:
diff --git a/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h
index 935304fe40..def825e529 100644
--- a/lib/Sema/TreeTransform.h
+++ b/lib/Sema/TreeTransform.h
@@ -26,6 +26,7 @@
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtObjC.h"
 #include "clang/AST/StmtOpenMP.h"
+#include "clang/AST/StmtPlatin.h"
 #include "clang/Sema/Designator.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Ownership.h"
@@ -1368,6 +1369,15 @@ public:
     return getSema().BuildObjCAtThrowStmt(AtLoc, Operand);
   }
 
+  /// \brief Build a new flow fact.
+  ///
+  StmtResult RebuildFlowfact(SourceRange Range) {
+    SmallVector<int, 2> foo;
+    SmallVector<std::string, 2> bar;
+    llvm_unreachable("not implemented");
+    return getSema().ActOnFlowfact(Range.getBegin(), Range.getEnd(), foo, bar, -1);
+  }
+
   /// \brief Build a new OpenMP executable directive.
   ///
   /// By default, performs semantic analysis to build the new statement.
@@ -11629,6 +11639,12 @@ TreeTransform<Derived>::TransformCapturedStmt(CapturedStmt *S) {
   return getSema().ActOnCapturedRegionEnd(Body.get());
 }
 
+template<typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformFlowfact(Flowfact *F) {
+  return getDerived().RebuildFlowfact(F->getSourceRange());
+}
+
 } // end namespace clang
 
 #endif // LLVM_CLANG_LIB_SEMA_TREETRANSFORM_H
diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp
index ad81ac8442..df74abcd0f 100644
--- a/lib/Serialization/ASTReaderStmt.cpp
+++ b/lib/Serialization/ASTReaderStmt.cpp
@@ -2476,6 +2476,10 @@ void ASTStmtReader::VisitOMPDistributeDirective(OMPDistributeDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitFlowfact(Flowfact *F) {
+  VisitStmt(F);
+}
+
 //===----------------------------------------------------------------------===//
 // ASTReader Implementation
 //===----------------------------------------------------------------------===//
@@ -2964,6 +2968,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
                                               nullptr);
       break;
 
+    case STMT_PLATIN:
+      llvm_unreachable("deserializing flowfact not implemented");
+      break;
+
     case STMT_OMP_PARALLEL_DIRECTIVE:
       S =
         OMPParallelDirective::CreateEmpty(Context,
diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp
index 000a2185f5..7a1596f5d7 100644
--- a/lib/Serialization/ASTWriterStmt.cpp
+++ b/lib/Serialization/ASTWriterStmt.cpp
@@ -2281,6 +2281,11 @@ void ASTStmtWriter::VisitOMPDistributeDirective(OMPDistributeDirective *D) {
   Code = serialization::STMT_OMP_DISTRIBUTE_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitFlowfact(Flowfact *F) {
+  VisitStmt(F);
+  Code = serialization::STMT_PLATIN;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
diff --git a/lib/Serialization/CMakeLists.txt b/lib/Serialization/CMakeLists.txt
index 95b33c388c..8aa9a7109a 100644
--- a/lib/Serialization/CMakeLists.txt
+++ b/lib/Serialization/CMakeLists.txt
@@ -12,6 +12,7 @@ add_clang_library(clangSerialization
   ASTWriter.cpp
   ASTWriterDecl.cpp
   ASTWriterStmt.cpp
+  FlowfactExporter.cpp
   GeneratePCH.cpp
   GlobalModuleIndex.cpp
   Module.cpp
diff --git a/lib/Serialization/FlowfactExporter.cpp b/lib/Serialization/FlowfactExporter.cpp
index be2424fb01..4e7cf266d4 100644
--- a/lib/Serialization/FlowfactExporter.cpp
+++ b/lib/Serialization/FlowfactExporter.cpp
@@ -90,13 +90,13 @@ namespace  {
 
       tool_output_file *OutFile;
       yaml::Output *Output;
-      std::string ErrorInfo;
+      std::error_code ErrorInfo;
 
-      OutFile = new tool_output_file(OutFileName.str().c_str(), ErrorInfo);
-      if (!ErrorInfo.empty()) {
+      OutFile = new tool_output_file(OutFileName, ErrorInfo, sys::fs::F_Text);
+      if (ErrorInfo) {
         delete OutFile;
         errs() << "[clang-ff] Opening Export File failed: " << OutFileName << "\n";
-        errs() << "[clang-ff] Reason: " << ErrorInfo;
+        errs() << "[clang-ff] Reason: " << ErrorInfo.message();
         return;
       }
       else {
@@ -113,6 +113,8 @@ namespace  {
   };
 }
 
-ASTConsumer *clang::CreateFlowfactExporter(StringRef filename) {
-  return new FlowfactExporter(filename);
+namespace clang {
+std::unique_ptr<ASTConsumer> CreateFlowfactExporter(StringRef filename) {
+  return llvm::make_unique<FlowfactExporter>(filename);
+}
 }
diff --git a/lib/StaticAnalyzer/Core/ExprEngine.cpp b/lib/StaticAnalyzer/Core/ExprEngine.cpp
index 662b0a2dd7..7b42a6c588 100644
--- a/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -808,6 +808,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
     case Stmt::WhileStmtClass:
     case Expr::MSDependentExistsStmtClass:
     case Stmt::CapturedStmtClass:
+    case Stmt::FlowfactClass:
     case Stmt::OMPParallelDirectiveClass:
     case Stmt::OMPSimdDirectiveClass:
     case Stmt::OMPForDirectiveClass:
diff --git a/test/CodeGenCUDA/address-space-conversion.cu b/test/CodeGenCUDA/address-space-conversion.cu
deleted file mode 100644
index ef88e2c7cd..0000000000
--- a/test/CodeGenCUDA/address-space-conversion.cu
+++ /dev/null
@@ -1,48 +0,0 @@
-// RUN: %clang_cc1 %s -triple nvptx-- -fcuda-is-device -emit-llvm -o - | FileCheck %s
-
-#include "../SemaCUDA/cuda.h"
-
-#define N 32
-
-extern __shared__ int x;
-
-
-__global__ void explicit_address_space_cast(int* p) {
-	// CHECK: @_Z27explicit_address_space_castPi
-   __shared__ unsigned char x[N];
-
-   for (unsigned int i=0; i<(N/4); i++) {
-     ((unsigned int *)x)[i] = 0;
-		// CHECK: addrspacecast
-   }
-}
-
-__global__ void pointer_as_array_access() {
-	// CHECK: @_Z23pointer_as_array_accessv
-   __shared__ int A[10];
-   int* p = A + 1;
-   p[x] = 0;
-	 // CHECK: addrspacecast
-}
-
-__device__ int* callee(int* p) {
-	// CHECK: @_Z6calleePi
-  return p;
-}
-
-__global__ void caller() {
-	// CHECK: @_Z6callerv
-  __shared__ int A[10];
-  __shared__ int* p;
-	p = A;
-	// CHECK: addrspacecast
-
-	((int*)A)[x] = 42;
-	// CHECK: addrspacecast
-	((int*)A)[0] = 15;
-	// CHECK: addrspacecast
-
-  int *np = callee(p);
-	A[2] = 5;
-	np[0] = 2;
-}
diff --git a/test/CodeGenOpenCL/address-space-conversion.cl b/test/CodeGenOpenCL/address-space-conversion.cl
deleted file mode 100644
index 6ae7dad520..0000000000
--- a/test/CodeGenOpenCL/address-space-conversion.cl
+++ /dev/null
@@ -1,14 +0,0 @@
-// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
-
-#define NULL ((void*)0)
-
-void null_pointer_implicit_conversion(int i, __global int *A) {
-	// CHECK: null_pointer_implicit_conversion
-	__global int *b;
-
-	b = i > 42 ? A : NULL;
-
-	if (b != NULL)
-	  A[0] = b[5];
-	// CHECK: null
-}
diff --git a/tools/libclang/CIndex.cpp b/tools/libclang/CIndex.cpp
index 9086c60e18..f40607091c 100644
--- a/tools/libclang/CIndex.cpp
+++ b/tools/libclang/CIndex.cpp
@@ -4466,6 +4466,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
     return cxstring::createRef("CXXAccessSpecifier");
   case CXCursor_ModuleImportDecl:
     return cxstring::createRef("ModuleImport");
+  case CXCursor_Flowfact:
+    return cxstring::createRef("Flowfact");
   case CXCursor_OMPParallelDirective:
     return cxstring::createRef("OMPParallelDirective");
   case CXCursor_OMPSimdDirective:
diff --git a/tools/libclang/CXCursor.cpp b/tools/libclang/CXCursor.cpp
index c766d2d69f..3689a15410 100644
--- a/tools/libclang/CXCursor.cpp
+++ b/tools/libclang/CXCursor.cpp
@@ -534,6 +534,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
   case Stmt::MSDependentExistsStmtClass:
     K = CXCursor_UnexposedStmt;
     break;
+  case Stmt::FlowfactClass:
+    K = CXCursor_Flowfact;
+    break;
   case Stmt::OMPParallelDirectiveClass:
     K = CXCursor_OMPParallelDirective;
     break;
-- 
GitLab