Adds AST matchers for matching CUDA declarations.

Patch by Jacques Pienaar.

llvm-svn: 214852
diff --git a/clang/unittests/ASTMatchers/ASTMatchersTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersTest.cpp
index e769e88..bf4dee8 100644
--- a/clang/unittests/ASTMatchers/ASTMatchersTest.cpp
+++ b/clang/unittests/ASTMatchers/ASTMatchersTest.cpp
@@ -649,6 +649,24 @@
   EXPECT_TRUE(matches("void f() { int i; }", CannotMemoize));
 }
 
+TEST(DeclarationMatcher, MatchCudaDecl) {
+  EXPECT_TRUE(matchesWithCuda("__global__ void f() { }"
+                              "void g() { f<<<1, 2>>>(); }",
+                              CUDAKernelCallExpr()));
+  EXPECT_TRUE(matchesWithCuda("__attribute__((device)) void f() {}",
+                              hasCudaDeviceAttr()));
+  EXPECT_TRUE(matchesWithCuda("__attribute__((host)) void f() {}",
+                              hasCudaHostAttr()));
+  EXPECT_TRUE(matchesWithCuda("__attribute__((global)) void f() {}",
+                              hasCudaGlobalAttr()));
+  EXPECT_FALSE(matchesWithCuda("void f() {}",
+                               hasCudaGlobalAttr()));
+  EXPECT_TRUE(notMatchesWithCuda("void f() {}",
+                                 hasCudaGlobalAttr()));
+  EXPECT_FALSE(notMatchesWithCuda("__attribute__((global)) void f() {}",
+                                  hasCudaGlobalAttr()));
+}
+
 // Implements a run method that returns whether BoundNodes contains a
 // Decl bound to Id that can be dynamically cast to T.
 // Optionally checks that the check succeeded a specific number of times.