Index: lib/Format/Format.cpp
===================================================================
--- lib/Format/Format.cpp
+++ lib/Format/Format.cpp
@@ -109,10 +109,8 @@
   }
 };
 
-template <>
-struct ScalarEnumerationTraits<FormatStyle::PointerAlignmentStyle> {
-  static void enumeration(IO &IO,
-                          FormatStyle::PointerAlignmentStyle &Value) {
+template <> struct ScalarEnumerationTraits<FormatStyle::PointerAlignmentStyle> {
+  static void enumeration(IO &IO, FormatStyle::PointerAlignmentStyle &Value) {
     IO.enumCase(Value, "Middle", FormatStyle::PAS_Middle);
     IO.enumCase(Value, "Left", FormatStyle::PAS_Left);
     IO.enumCase(Value, "Right", FormatStyle::PAS_Right);
@@ -144,8 +142,8 @@
     IO.mapOptional("Language", Style.Language);
 
     if (IO.outputting()) {
-      StringRef StylesArray[] = { "LLVM",    "Google", "Chromium",
-                                  "Mozilla", "WebKit", "GNU" };
+      StringRef StylesArray[] = {"LLVM",    "Google", "Chromium",
+                                 "Mozilla", "WebKit", "GNU"};
       ArrayRef<StringRef> Styles(StylesArray);
       for (size_t i = 0, e = Styles.size(); i < e; ++i) {
         StringRef StyleName(Styles[i]);
@@ -273,7 +271,7 @@
 // will be used to get default values for missing keys.
 // If the first element has no Language specified, it will be treated as the
 // default one for the following elements.
-template <> struct DocumentListTraits<std::vector<FormatStyle> > {
+template <> struct DocumentListTraits<std::vector<FormatStyle>> {
   static size_t size(IO &IO, std::vector<FormatStyle> &Seq) {
     return Seq.size();
   }
@@ -632,6 +630,10 @@
     if (tryMergeConflictMarkers())
       return;
 
+    if (Style.Language == FormatStyle::LK_Cpp)
+      if (tryMergeCudaTokens())
+        return;
+
     if (Style.Language == FormatStyle::LK_JavaScript) {
       if (tryMergeJSRegexLiteral())
         return;
@@ -638,11 +640,11 @@
       if (tryMergeEscapeSequence())
         return;
 
-      static tok::TokenKind JSIdentity[] = { tok::equalequal, tok::equal };
-      static tok::TokenKind JSNotIdentity[] = { tok::exclaimequal, tok::equal };
-      static tok::TokenKind JSShiftEqual[] = { tok::greater, tok::greater,
-                                               tok::greaterequal };
-      static tok::TokenKind JSRightArrow[] = { tok::equal, tok::greater };
+      static tok::TokenKind JSIdentity[] = {tok::equalequal, tok::equal};
+      static tok::TokenKind JSNotIdentity[] = {tok::exclaimequal, tok::equal};
+      static tok::TokenKind JSShiftEqual[] = {tok::greater, tok::greater,
+                                              tok::greaterequal};
+      static tok::TokenKind JSRightArrow[] = {tok::equal, tok::greater};
       // FIXME: We probably need to change token type to mimic operator with the
       // correct priority.
       if (tryMergeTokens(JSIdentity))
@@ -666,8 +668,9 @@
       return false;
     unsigned AddLength = 0;
     for (unsigned i = 1; i < Kinds.size(); ++i) {
-      if (!First[i]->is(Kinds[i]) || First[i]->WhitespaceRange.getBegin() !=
-                                         First[i]->WhitespaceRange.getEnd())
+      if (!First[i]->is(Kinds[i]) ||
+          First[i]->WhitespaceRange.getBegin() !=
+              First[i]->WhitespaceRange.getEnd())
         return false;
       AddLength += First[i]->TokenText.size();
     }
@@ -678,6 +681,42 @@
     return true;
   }
 
+  bool tryMergeCudaTokens() {
+    static tok::TokenKind TrippleOpen[] = {tok::lessless, tok::less};
+    static tok::TokenKind TrippleClose[] = {tok::greater, tok::greater,
+                                            tok::greater};
+    // Merge lessless, less into lesslessless
+    if (tryMergeTokens(TrippleOpen)) {
+      SmallVectorImpl<FormatToken *>::const_iterator Merged = Tokens.end() - 1;
+      (*Merged)->Tok.setKind(tok::lesslessless);
+      return true;
+    }
+
+    // Only merge greater,greater,greater into greatergreatergreater if an
+    // lesslessless was seen before the last semi
+    if (Tokens.size() < 3)
+      return false;
+    if (Tokens[Tokens.size() - 1]->is(tok::greater) &&
+        Tokens[Tokens.size() - 2]->is(tok::greater) &&
+        Tokens[Tokens.size() - 3]->is(tok::greater)) {
+      for (auto I = Tokens.rbegin() + 1, E = Tokens.rend(); I != E; ++I) {
+        if ((*I)->is(tok::semi))
+          return false;
+        if ((*I)->is(tok::lesslessless)) {
+          if (tryMergeTokens(TrippleClose)) {
+            SmallVectorImpl<FormatToken *>::const_iterator Merged =
+                Tokens.end() - 1;
+            (*Merged)->Tok.setKind(tok::greatergreatergreater);
+            return true;
+          }
+          return false;
+        }
+      }
+    }
+
+    return false;
+  }
+
   // Tries to merge an escape sequence, i.e. a "\\" and the following
   // character. Use e.g. inside JavaScript regex literals.
   bool tryMergeEscapeSequence() {
Index: lib/Format/TokenAnnotator.cpp
===================================================================
--- lib/Format/TokenAnnotator.cpp
+++ lib/Format/TokenAnnotator.cpp
@@ -1807,6 +1807,18 @@
              Left.isOneOf(TT_TemplateCloser, TT_TemplateOpener));
   if ((Left.is(TT_TemplateOpener)) != (Right.is(TT_TemplateCloser)))
     return Style.SpacesInAngles;
+  // Treat parameters in between <<< and >>> similar to template parameters
+  if (Left.is(tok::lesslessless) && Right.isNot(tok::greatergreatergreater))
+    return Style.SpacesInAngles;
+  if (Right.is(tok::greatergreatergreater))
+    return Style.SpacesInAngles;
+  // No space needed between identifier or ">" and "<<<".
+  if ((Left.is(tok::identifier) || Left.is(tok::greater)) &&
+      Right.is(tok::lesslessless))
+    return false;
+  // No space needed between ">>>" and "("
+  if (Right.is(tok::l_paren) && Left.is(tok::greatergreatergreater))
+    return false;
   if ((Right.is(TT_BinaryOperator) && !Left.is(tok::l_paren)) ||
       Left.isOneOf(TT_BinaryOperator, TT_ConditionalExpr))
     return true;
Index: unittests/Format/FormatTest.cpp
===================================================================
--- unittests/Format/FormatTest.cpp
+++ unittests/Format/FormatTest.cpp
@@ -9646,6 +9646,18 @@
   verifyFormat("A<A<int>>();", Spaces);
 }
 
+TEST_F(FormatTest, TripleAngleBrackets) {
+  verifyFormat("f<<<1, 1>>>();");
+  verifyFormat("f<<<1, 1, 1, s>>>();");
+  verifyFormat("f<<<a, b, c, d>>>();");
+  EXPECT_EQ("f<<<1, 1>>>();",
+            format("f <<< 1, 1 >>> ();"));
+  verifyFormat("f<param><<<1, 1>>>();");
+  verifyFormat("f<1><<<1, 1>>>();");
+  EXPECT_EQ("f<param><<<1, 1>>>();",
+            format("f< param > <<< 1, 1 >>> ();"));
+}
+
 TEST_F(FormatTest, HandleUnbalancedImplicitBracesAcrossPPBranches) {
   std::string code = "#if A\n"
                      "#if B\n"