Index: docs/LibASTMatchersReference.html
===================================================================
--- docs/LibASTMatchersReference.html
+++ docs/LibASTMatchersReference.html
@@ -1612,6 +1612,16 @@
 </pre></td></tr>
 
 
+<tr><td>Matcher&lt<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>&gt;</td><td class="name" onclick="toggle('hasAttr0')"><a name="hasAttr0Anchor">hasAttr</a></td><td></td></tr>
+<tr><td colspan="4" class="doc" id="hasAttr0"><pre>Matches declaration that has a given attribute.
+
+Given
+  __attribute__((device)) void f() { ... }
+decl(hasAttr&lt;clang::CUDADeviceAttr&gt;()) matches the function declaration of
+f.
+</pre></td></tr>
+
+
 <tr><td>Matcher&lt<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>&gt;</td><td class="name" onclick="toggle('hasCudaDeviceAttr0')"><a name="hasCudaDeviceAttr0Anchor">hasCudaDeviceAttr</a></td><td></td></tr>
 <tr><td colspan="4" class="doc" id="hasCudaDeviceAttr0"><pre>Matches declaration that has CUDA device attribute.
 
Index: include/clang/ASTMatchers/ASTMatchers.h
===================================================================
--- include/clang/ASTMatchers/ASTMatchers.h
+++ include/clang/ASTMatchers/ASTMatchers.h
@@ -3658,6 +3658,19 @@
   return InnerMatcher.matches(*Node.getLHS(), Finder, Builder);
 }
 
+/// \brief Matches declaration that has a given attribute.
+///
+/// Given
+/// \code
+///   __attribute__((device)) void f() { ... }
+/// \endcode
+/// decl(hasAttr<clang::CUDADeviceAttr>()) matches the function declaration of
+/// f.
+template <typename Attr>
+inline internal::Matcher<Decl> hasAttr() {
+  return internal::HasAttrMatcher<Decl, Attr>();
+}
+
 /// \brief Matches CUDA kernel call expression.
 ///
 /// Example matches,
@@ -3674,8 +3687,8 @@
 ///   __attribute__((device)) void f() { ... }
 /// \endcode
 /// matches the function declaration of f.
-AST_MATCHER(Decl, hasCudaDeviceAttr) {
-  return Node.hasAttr<clang::CUDADeviceAttr>();
+inline internal::Matcher<Decl> hasCudaDeviceAttr() {
+  return hasAttr<CUDADeviceAttr>();
 }
 
 /// \brief Matches declaration that has CUDA host attribute.
@@ -3685,8 +3698,8 @@
 ///   __attribute__((host)) void f() { ... }
 /// \endcode
 /// matches the function declaration of f.
-AST_MATCHER(Decl, hasCudaHostAttr) {
-  return Node.hasAttr<clang::CUDAHostAttr>();
+inline internal::Matcher<Decl> hasCudaHostAttr() {
+  return hasAttr<CUDAHostAttr>();
 }
 
 /// \brief  Matches declaration that has CUDA global attribute.
@@ -3696,8 +3709,8 @@
 ///   __attribute__((global)) void f() { ... }
 /// \endcode
 /// matches the function declaration of f.
-AST_MATCHER(Decl, hasCudaGlobalAttr) {
-  return Node.hasAttr<clang::CUDAGlobalAttr>();
+inline internal::Matcher<Decl> hasCudaGlobalAttr() {
+  return hasAttr<CUDAGlobalAttr>();
 }
 
 } // end namespace ast_matchers
Index: include/clang/ASTMatchers/ASTMatchersInternal.h
===================================================================
--- include/clang/ASTMatchers/ASTMatchersInternal.h
+++ include/clang/ASTMatchers/ASTMatchersInternal.h
@@ -1658,6 +1658,28 @@
   ast_type_traits::DynTypedNode Node;
 };
 
+/// \brief Matches declarations have an attribute of type \c Attr
+template <typename T, typename Attr,
+          typename ReturnTypesF = void(internal::TypeList<Decl>)>
+class HasAttrMatcher : public SingleNodeMatcherInterface<T> {
+  static_assert(std::is_same<T, Decl>::value,
+                "unsupported class for matcher");
+public:
+  typedef typename ExtractFunctionArgMeta<ReturnTypesF>::type ReturnTypes;
+  operator Matcher<T>() const {
+    static_assert(TypeListContainsSuperOf<ReturnTypes, T>::value,
+                  "right polymorphic conversion");
+    return Matcher<T>(new HasAttrMatcher<T, Attr, ReturnTypesF>());
+  }
+  bool matchesNode(const T &Node) const override {
+    return matchesSpecialized(Node);
+  }
+private:
+  bool matchesSpecialized(const Decl &Node) const {
+    return Node.hasAttr<Attr>();
+  }
+};
+
 } // end namespace internal
 } // end namespace ast_matchers
 } // end namespace clang
Index: lib/ASTMatchers/Dynamic/Registry.cpp
===================================================================
--- lib/ASTMatchers/Dynamic/Registry.cpp
+++ lib/ASTMatchers/Dynamic/Registry.cpp
@@ -78,6 +78,7 @@
   //
   // Polymorphic + argument overload:
   // findAll
+  // hasAttr
   //
   // Other:
   // equals
Index: unittests/ASTMatchers/ASTMatchersTest.cpp
===================================================================
--- unittests/ASTMatchers/ASTMatchersTest.cpp
+++ unittests/ASTMatchers/ASTMatchersTest.cpp
@@ -649,6 +649,13 @@
   EXPECT_TRUE(matches("void f() { int i; }", CannotMemoize));
 }
 
+TEST(DeclarationMatcher, HasAttr) {
+  EXPECT_TRUE(matches("struct __attribute__((warn_unused)) X {};",
+                      decl(hasAttr<WarnUnusedAttr>())));
+  EXPECT_FALSE(matches("struct X {};",
+                       decl(hasAttr<WarnUnusedAttr>())));
+}
+
 TEST(DeclarationMatcher, MatchCudaDecl) {
   EXPECT_TRUE(matchesWithCuda("__global__ void f() { }"
                               "void g() { f<<<1, 2>>>(); }",