diff --git a/clang/docs/LibASTMatchersReference.html b/clang/docs/LibASTMatchersReference.html --- a/clang/docs/LibASTMatchersReference.html +++ b/clang/docs/LibASTMatchersReference.html @@ -6426,6 +6426,46 @@ +Matcher<CUDAKernelCallExpr>cudaBlockDimMatcher<Expr> InnerMatcher +
Matches the second argument (block dim) in <<<>>> on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaBlockDim()) will match j in
+  kernel<<<i,j>>>();
+
+ + +Matcher<CUDAKernelCallExpr>cudaGridDimMatcher<Expr> InnerMatcher +
Matches the first argument (grid dim) in <<<>>> on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaGridDim()) will match i in
+  kernel<<<i,j>>>();
+
+ + +Matcher<CUDAKernelCallExpr>cudaSharedMemPerBlockMatcher<Expr> InnerMatcher +
Matches the third argument (shared mem size) in <<<>>> on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaSharedMemPerBlock()) will match mem in
+  kernel<<<i, j, mem, 0>>>();
+
+ + +Matcher<CUDAKernelCallExpr>cudaStreamMatcher<Expr> InnerMatcher +
Matches the fourth argument (CUDA stream) in <<<>>> on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaStream()) will match 0 in
+  kernel<<<i, j, mem, 0>>>();
+
+ + +Matcher<CUDAKernelCallExpr>hasKernelConfigMatcher<CallExpr> InnerMatcher +
Matches the kernel launch config (in <<<>>>) on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(hasKernelConfig()) will match <<<i,j>>> in
+  kernel<<<i,j>>>();
+
+ + Matcher<CXXBaseSpecifier>hasTypeLocMatcher<TypeLoc> Inner 
Matches if the type location of a node matches the inner matcher.
 
diff --git a/clang/include/clang/ASTMatchers/ASTMatchers.h b/clang/include/clang/ASTMatchers/ASTMatchers.h
--- a/clang/include/clang/ASTMatchers/ASTMatchers.h
+++ b/clang/include/clang/ASTMatchers/ASTMatchers.h
@@ -7848,6 +7848,80 @@
 extern const internal::VariadicDynCastAllOfMatcher
     cudaKernelCallExpr;
 
+/// Matches the kernel launch config (in <<<>>>) on CUDA kernel calls.
+///
+/// Example: cudaKernelCallExpr(hasKernelConfig()) will match <<>> in
+/// \code
+///   kernel<<>>();
+/// \endcode
+AST_MATCHER_P(CUDAKernelCallExpr, hasKernelConfig, internal::Matcher,
+              InnerMatcher) {
+  if (const CallExpr *Config = Node.getConfig()) {
+    return InnerMatcher.matches(*Config, Finder, Builder);
+  }
+  return false;
+}
+
+/// Matches the first argument (grid dim) in <<<>>> on CUDA kernel calls.
+///
+/// Example: cudaKernelCallExpr(cudaGridDim()) will match i in
+/// \code
+///   kernel<<>>();
+/// \endcode
+AST_MATCHER_P(CUDAKernelCallExpr, cudaGridDim, internal::Matcher,
+              InnerMatcher) {
+  const CallExpr *Config = Node.getConfig();
+  if (Config && Config->getNumArgs() > 0) {
+    return InnerMatcher.matches(*(Config->getArg(0)), Finder, Builder);
+  }
+  return false;
+}
+
+/// Matches the second argument (block dim) in <<<>>> on CUDA kernel calls.
+///
+/// Example: cudaKernelCallExpr(cudaBlockDim()) will match j in
+/// \code
+///   kernel<<>>();
+/// \endcode
+AST_MATCHER_P(CUDAKernelCallExpr, cudaBlockDim, internal::Matcher,
+              InnerMatcher) {
+  const CallExpr *Config = Node.getConfig();
+  if (Config && Config->getNumArgs() > 1) {
+    return InnerMatcher.matches(*(Config->getArg(1)), Finder, Builder);
+  }
+  return false;
+}
+
+/// Matches the third argument (shared mem size) in <<<>>> on CUDA kernel calls.
+///
+/// Example: cudaKernelCallExpr(cudaSharedMemPerBlock()) will match mem in
+/// \code
+///   kernel<<>>();
+/// \endcode
+AST_MATCHER_P(CUDAKernelCallExpr, cudaSharedMemPerBlock,
+              internal::Matcher, InnerMatcher) {
+  const CallExpr *Config = Node.getConfig();
+  if (Config && Config->getNumArgs() > 2) {
+    return InnerMatcher.matches(*(Config->getArg(2)), Finder, Builder);
+  }
+  return false;
+}
+
+/// Matches the fourth argument (CUDA stream) in <<<>>> on CUDA kernel calls.
+///
+/// Example: cudaKernelCallExpr(cudaStream()) will match 0 in
+/// \code
+///   kernel<<>>();
+/// \endcode
+AST_MATCHER_P(CUDAKernelCallExpr, cudaStream, internal::Matcher,
+              InnerMatcher) {
+  const CallExpr *Config = Node.getConfig();
+  if (Config && Config->getNumArgs() > 3) {
+    return InnerMatcher.matches(*(Config->getArg(3)), Finder, Builder);
+  }
+  return false;
+}
+
 /// Matches expressions that resolve to a null pointer constant, such as
 /// GNU's __null, C++11's nullptr, or C's NULL macro.
 ///
diff --git a/clang/lib/ASTMatchers/Dynamic/Registry.cpp b/clang/lib/ASTMatchers/Dynamic/Registry.cpp
--- a/clang/lib/ASTMatchers/Dynamic/Registry.cpp
+++ b/clang/lib/ASTMatchers/Dynamic/Registry.cpp
@@ -177,7 +177,11 @@
   REGISTER_MATCHER(continueStmt);
   REGISTER_MATCHER(coreturnStmt);
   REGISTER_MATCHER(coyieldExpr);
+  REGISTER_MATCHER(cudaBlockDim);
+  REGISTER_MATCHER(cudaGridDim);
   REGISTER_MATCHER(cudaKernelCallExpr);
+  REGISTER_MATCHER(cudaSharedMemPerBlock);
+  REGISTER_MATCHER(cudaStream);
   REGISTER_MATCHER(cxxBaseSpecifier);
   REGISTER_MATCHER(cxxBindTemporaryExpr);
   REGISTER_MATCHER(cxxBoolLiteral);
@@ -320,6 +324,7 @@
   REGISTER_MATCHER(hasInit);
   REGISTER_MATCHER(hasInitializer);
   REGISTER_MATCHER(hasInitStatement);
+  REGISTER_MATCHER(hasKernelConfig);
   REGISTER_MATCHER(hasKeywordSelector);
   REGISTER_MATCHER(hasLHS);
   REGISTER_MATCHER(hasLocalQualifiers);